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

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.

Reviewers: rnk, echristo

Subscribers: cfe-commits, jhen, tra

Differential Revision: http://reviews.llvm.org/D16664

llvm-svn: 259122

1 files changed, 23 insertions, 33 deletions

diff --git a/clang/test/CodeGenCUDA/printf.cu b/clang/test/CodeGenCUDA/printf.cu
index f91aba78784..dc3f4ea788f 100644
--- a/clang/test/CodeGenCUDA/printf.cu
+++ b/clang/test/CodeGenCUDA/printf.cu
@@ -9,45 +9,35 @@
 extern "C" __device__ int vprintf(const char*, const char*);
 
 // Check a simple call to printf end-to-end.
+// CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double }
 __device__ int CheckSimple() {
+  // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]
   // 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]])
+  const char* fmt = "%d %lld %f";
+  // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 0
+  // CHECK: store i32 1, i32* [[PTR0]], align 4
+  // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 1
+  // CHECK: store i64 2, i64* [[PTR1]], align 8
+  // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds [[SIMPLE_PRINTF_TY]], [[SIMPLE_PRINTF_TY]]* [[BUF]], i32 0, i32 2
+  // CHECK: store double 3.0{{[^,]*}}, double* [[PTR2]], align 8
+  // CHECK: [[BUF_CAST:%[0-9]+]] = bitcast [[SIMPLE_PRINTF_TY]]* [[BUF]] to i8*
+  // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF_CAST]])
   // CHECK: ret i32 [[RET]]
-  return printf(fmt, 42);
-}
-
-// Check that the args' types are promoted correctly when we call printf.
-__device__ void CheckTypes() {
-  // CHECK: alloca {{.*}} align 8
-  // CHECK: getelementptr {{.*}} i32 0
-  // CHECK: bitcast {{.*}} to i32*
-  // CHECK: getelementptr {{.*}} i32 4
-  // CHECK: bitcast {{.*}} to i32*
-  // CHECK: getelementptr {{.*}} i32 8
-  // CHECK: bitcast {{.*}} to double*
-  // CHECK: getelementptr {{.*}} i32 16
-  // CHECK: bitcast {{.*}} to double*
-  printf("%d %d %f %f", (char)1, (short)2, 3.0f, 4.0);
-}
-
-// Check that the args are aligned properly in the buffer.
-__device__ void CheckAlign() {
-  // CHECK: alloca i8, i32 40, align 8
-  // CHECK: getelementptr {{.*}} i32 0
-  // CHECK: getelementptr {{.*}} i32 8
-  // CHECK: getelementptr {{.*}} i32 16
-  // CHECK: getelementptr {{.*}} i32 20
-  // CHECK: getelementptr {{.*}} i32 24
-  // CHECK: getelementptr {{.*}} i32 32
-  printf("%d %f %d %d %d %lld", 1, 2.0, 3, 4, 5, (long long)6);
+  return printf(fmt, 1, 2ll, 3.0);
 }
 
 __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 %printf_args
+  // CHECK: call {{.*}} @_Z3foov()
+  if (foo()) {
+    printf("%d", 42);
+  }
+}