[CUDA] Make printf work.

Summary:
The code in CGCUDACall is largely based on a patch written by Eli
Bendersky:
http://lists.llvm.org/pipermail/llvm-commits/Week-of-Mon-20140324/210218.html

That patch implemented an LLVM pass lowering printf to vprintf; this
one does something similar, but in Clang codegen.

Reviewers: echristo

Subscribers: cfe-commits, jhen, tra, majnemer

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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@258642 91177308-0d34-0410-b5e6-96231b3b80d8

2 files changed, 55 insertions, 0 deletions

diff --git a/test/CodeGenCUDA/Inputs/cuda.h b/test/CodeGenCUDA/Inputs/cuda.h
index a9a4595a14..9b9f43a1aa 100644
--- a/test/CodeGenCUDA/Inputs/cuda.h
+++ b/test/CodeGenCUDA/Inputs/cuda.h
@@ -18,3 +18,5 @@ typedef struct cudaStream *cudaStream_t;
 
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
+
+extern "C" __device__ int printf(const char*, ...);
diff --git a/test/CodeGenCUDA/printf.cu b/test/CodeGenCUDA/printf.cu
new file mode 100644
index 0000000000..f91aba7878
--- /dev/null
+++ b/test/CodeGenCUDA/printf.cu
@@ -0,0 +1,53 @@
+// 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: [[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);
+}
+
+// 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);
+}
+
+__device__ void CheckNoArgs() {
+  // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}}
+  printf("hello, world!");
+}