[CUDA] Generate CUDA's printf alloca in its function's entry block.

[jlebar]

Summary

This is necessary to prevent llvm from generating stacksave intrinsics
around this alloca. NVVM doesn't have a stack, and we don't handle said
intrinsics.

[echristo]

lib/CodeGen/CGCUDABuiltin.cpp

109

Not quite, you'll want to use AllocaInsertPt for this or even CreateTempAlloca.

[jlebar]

Address echristo's review comments.

[jlebar]

lib/CodeGen/CGCUDABuiltin.cpp

109

Aha. Used AllocaInsertPt because it doesn't seem that there's an overload of CreateTempAlloca that takes an explicit size.

[rnk]

lib/CodeGen/CGCUDABuiltin.cpp

105–108

The fact that allocas for local variables should always go in the entry block is pretty widespread cultural knowledge in LLVM and clang. Most readers aren't going to need this comment, unless you expect that people working on CUDA won't have that background. Plus, if you use CreateTempAlloca, there won't be any question about which insert point should be used.

109

You can still use CreateTempAlloca by making an [i8 x N] LLVM type. You'll have to use CreateStructGEP below for forming GEPs. Overall I think that'd be nicer, since you don't need to worry about insertion at all.

[echristo]

lib/CodeGen/CGCUDABuiltin.cpp

109

+1 :)

112

Also you'd have wanted to insert it before anyhow.

[jlebar]

Use a struct rather than an i8 buffer.

[jlebar]

Thank you for the reviews.

Please have another look; I switched to using a struct proper. It's a lot cleaner! We're now assuming that the struct is aligned in the same way as vprintf wants, but if anything I expect this new code is more likely to match what it wants.

lib/CodeGen/CGCUDABuiltin.cpp

105–108

OK, yeah, I also don't like comments that explain something that everyone other than the author knows. Thanks.

[echristo]

One inline nit then LGTM.

-eric

lib/CodeGen/CGCUDABuiltin.cpp

102

• on the wrong side ;)

[jlebar]

lib/CodeGen/CGCUDABuiltin.cpp

102

Argh, I really need to set up a linter. I'm still doing readability reviews, and I cannot brain two styles. Sorry to keep wasting your time with silly stuff like this.

[echristo]

lib/CodeGen/CGCUDABuiltin.cpp

102

You could just use clang-format on everything :)

[jlebar]

Do you have a script that will take as input a commit range and git
commit --amend clang-tidy fixes for lines modified in that range?
Because if so,

a) I would be your best friend forever, and
b) It should be simple to convert that into a linter for arc to catch
the case when I forget to run said tool.

[jlebar]

Hm, well, https://llvm.org/svn/llvm-project/cfe/trunk/tools/clang-format/git-clang-format
is close... Not sure if that triggers the bff clause, will consult my
attorney.

lib/CodeGen/CGCUDABuiltin.cpp

for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I) {
     // Add padding required to keep the current arg aligned.
     BufSize = llvm::alignTo(BufSize, Align);
     BufSize += DL.getTypeAllocSize(Ty);
   }
 
   // Construct and fill the buffer.
   llvm::Value* BufferPtr = nullptr;
   if (BufSize == 0) {
     // If there are no args, pass a null pointer to vprintf.
     BufferPtr = llvm::ConstantPointerNull::get(llvm::Type::getInt8PtrTy(Ctx));
   } else {
-    BufferPtr = Builder.Insert(new llvm::AllocaInst(
+    // Insert our alloca not into the current BB, but into the function's entry
+    // block.  This is important because nvvm doesn't support alloca -- if we
+    // put the alloca anywhere else, llvm may eventually output
+    // stacksave/stackrestore intrinsics, which cause our nvvm backend to choke.
+    auto *Alloca = new llvm::AllocaInst(
         llvm::Type::getInt8Ty(Ctx), llvm::ConstantInt::get(Int32Ty, BufSize),
-        BufAlign, "printf_arg_buf"));
+        BufAlign, "printf_arg_buf");
+    Alloca->insertAfter(AllocaInsertPt);
+    BufferPtr = Alloca;
 
     unsigned Offset = 0;
     for (unsigned I = 1, NumArgs = Args.size(); I < NumArgs; ++I) {
       llvm::Value *Arg = Args[I].RV.getScalarVal();
       llvm::Type *Ty = Arg->getType();
       auto Align = DL.getPrefTypeAlignment(Ty);
 
       // Pad the buffer to Arg's alignment.

test/CodeGenCUDA/printf.cu

 // REQUIRES: x86-registered-target
 // REQUIRES: nvptx-registered-target
 
 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \
 // RUN:   -o - %s | FileCheck %s
 
 #include "Inputs/cuda.h"
 
 extern "C" __device__ int vprintf(const char*, const char*);
 
 // Check a simple call to printf end-to-end.
 __device__ int CheckSimple() {
+  // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca i8, i32 4, align 4
   // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt
   const char* fmt = "%d";
-  // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca i8, i32 4, align 4
   // CHECK: [[PTR:%[0-9]+]] = getelementptr i8, i8* [[BUF]], i32 0
   // CHECK: [[CAST:%[0-9]+]] = bitcast i8* [[PTR]] to i32*
   // CHECK: store i32 42, i32* [[CAST]], align 4
   // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF]])
   // CHECK: ret i32 [[RET]]
   return printf(fmt, 42);
 }
 
__device__ void CheckAlign() {
   // CHECK: getelementptr {{.*}} i32 32
   printf("%d %f %d %d %d %lld", 1, 2.0, 3, 4, 5, (long long)6);
 }
 
 __device__ void CheckNoArgs() {
   // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}}
   printf("hello, world!");
 }
+
+// Check that printf's alloca happens in the entry block, not inside the if
+// statement.
+__device__ bool foo();
+__device__ void CheckAllocaIsInEntryBlock() {
+  // CHECK: alloca i8, i32 4, align 4
+  // CHECK: call {{.*}} @_Z3foov()
+  if (foo()) {
+    printf("%d", 42);
+  }
+}