[NVPTX] CUDA does provide malloc/free since compute capability 2.X

[jdoerfert]

Summary

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-memory-allocation-and-operations

[tra]

Tests? NVPTX used to have issues lowering libcalls at all. It would be great to verify that it works now.

[jdoerfert]

In D98606#2626735, @tra wrote:

Tests? NVPTX used to have issues lowering libcalls at all. It would be great to verify that it works now.

We use malloc in the OpenMP runtime, it works, albeit slow. Where do I put a test for this?

[tra]

In D98606#2627001, @jdoerfert wrote:

In D98606#2626735, @tra wrote:

Tests? NVPTX used to have issues lowering libcalls at all. It would be great to verify that it works now.

We use malloc in the OpenMP runtime, it works, albeit slow. Where do I put a test for this?

test/CodeGen/NVPTX/libcall-fulfilled.ll might be a good place.

[jdoerfert]

Add tests, merge D98607 into this as comment.
TBH. This patch doesn't impact the CodeGen test but only the instcombine test.

[tra]

LGTM module couple of test nits.

llvm/test/CodeGen/NVPTX/libcall-fulfilled.ll

39–45

I'd just leave

CHECK: call.uni
CHECK: malloc

CHECK: call-uni
CHECK: free

The surrounding call setup machinery is not particularly interesting and can potentially morph due to unrelated backend changes.

llvm/test/Transforms/InstCombine/malloc_free_delete_nvptx.ll

13

I'd explicitly check that malloc/free are missing to make it obvious.

CHECK-NOT: malloc
CHECK-NOT: free
CHECK: ret void

Otherwise the test looks somewhat odd -- it's supposedly about mallof and free, but the CHECK lines don't mention either. The optimize them properly in the comment above could also be more specific about what we expect to happen.

[jdoerfert]

llvm/test/Transforms/InstCombine/malloc_free_delete_nvptx.ll

13

would it be ok to just improve the comment? CHECK-NEXT does the heavy lifting here and we might want to keep the auto-generated check lines.

[tra]

llvm/test/Transforms/InstCombine/malloc_free_delete_nvptx.ll

13

That'd work too.

[jdoerfert]

Addressed comments

llvm/lib/Analysis/TargetLibraryInfo.cpp

static void initialize(TargetLibraryInfoImpl &TLI, const Triple &T,
   // but only used functions are provided to llvm.  Moreover, most of the
   // functions in libdevice don't map precisely to standard library functions.
   //
   // FIXME: Having no standard library prevents e.g. many fastmath
   // optimizations, so this situation should be fixed.
   if (T.isNVPTX()) {
     TLI.disableAllFunctions();
     TLI.setAvailable(LibFunc_nvvm_reflect);
+    TLI.setAvailable(llvm::LibFunc_malloc);
+    TLI.setAvailable(llvm::LibFunc_free);
+
+    // TODO: We could enable the following two according to [0] but we haven't
+    //       done an evaluation wrt. the performance implications.
+    // [0]
+    // https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-memory-allocation-and-operations
+    //
+    //    TLI.setAvailable(llvm::LibFunc_memcpy);
+    //    TLI.setAvailable(llvm::LibFunc_memset);
+
   } else {
     TLI.setUnavailable(LibFunc_nvvm_reflect);
   }
 
   // These vec_malloc/free routines are only available on AIX.
   if (!T.isOSAIX()) {
     TLI.setUnavailable(LibFunc_vec_calloc);
     TLI.setUnavailable(LibFunc_vec_malloc);

llvm/test/CodeGen/NVPTX/libcall-fulfilled.ll

bb1:
   ret i128 %a
 }
 
 ; Underlying libcall definition
 ; CHECK: .visible .func  (.param .align 16 .b8 func_retval0[16]) __umodti3(
 define i128 @__umodti3(i128, i128) {
   ret i128 0
 }
+
+declare i8* @malloc(i64)
+declare void @free(i8*)
+
+define void @malloc_then_free() {
+; CHECK:  call.uni (retval0),
+; CHECK:  malloc,
+; CHECK:  (
+; CHECK:  param0
+; CHECK:  );
+; CHECK:  ld.param.b32 %r1, [retval0+0];
+; CHECK:  } // callseq 1
+; CHECK:  mov.u16 %rs1, 0;
+; CHECK:  st.u8 [%r1], %rs1;
+; CHECK:  { // callseq 2, 0
+; CHECK:  .reg .b32 temp_param_reg;
+; CHECK:  .param .b32 param0;
+; CHECK:  st.param.b32 [param0+0], %r1;
+; CHECK:  call.uni
+; CHECK:  free,
+; CHECK:  (
+; CHECK:  param0
+; CHECK:  );
+; CHECK:  } // callseq 2
+  %a = call i8* @malloc(i64 4)
+  store i8 0, i8* %a
+  call void @free(i8* %a)
+  ret void
+}

llvm/test/Transforms/InstCombine/malloc_free_delete_nvptx.ll

+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt < %s -instcombine -S | FileCheck %s
+
+; Ensure the nvptx backend states malloc & free are a thing so we can recognize 
+; and optimize them properly.
+target triple = "nvptx64"
+
+declare i8* @malloc(i64)
+declare void @free(i8*)
+
+define void @malloc_then_free() {
+; CHECK-LABEL: @malloc_then_free(
+; CHECK-NEXT:    ret void
+;
+  %a = call i8* @malloc(i64 4)
+  store i8 0, i8* %a
+  call void @free(i8* %a)
+  ret void
+}
+