1 // REQUIRES: x86-registered-target 2 // REQUIRES: nvptx-registered-target 3 4 // RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \ 5 // RUN: -o - %s | FileCheck %s 6 7 #include "Inputs/cuda.h" 8 9 extern "C" __device__ int vprintf(const char*, const char*); 10 11 // Check a simple call to printf end-to-end. 12 __device__ int CheckSimple() { 13 // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt 14 const char* fmt = "%d"; 15 // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca i8, i32 4, align 4 16 // CHECK: [[PTR:%[0-9]+]] = getelementptr i8, i8* [[BUF]], i32 0 17 // CHECK: [[CAST:%[0-9]+]] = bitcast i8* [[PTR]] to i32* 18 // CHECK: store i32 42, i32* [[CAST]], align 4 19 // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(i8* [[FMT]], i8* [[BUF]]) 20 // CHECK: ret i32 [[RET]] 21 return printf(fmt, 42); 22 } 23 24 // Check that the args' types are promoted correctly when we call printf. 25 __device__ void CheckTypes() { 26 // CHECK: alloca {{.*}} align 8 27 // CHECK: getelementptr {{.*}} i32 0 28 // CHECK: bitcast {{.*}} to i32* 29 // CHECK: getelementptr {{.*}} i32 4 30 // CHECK: bitcast {{.*}} to i32* 31 // CHECK: getelementptr {{.*}} i32 8 32 // CHECK: bitcast {{.*}} to double* 33 // CHECK: getelementptr {{.*}} i32 16 34 // CHECK: bitcast {{.*}} to double* 35 printf("%d %d %f %f", (char)1, (short)2, 3.0f, 4.0); 36 } 37 38 // Check that the args are aligned properly in the buffer. 39 __device__ void CheckAlign() { 40 // CHECK: alloca i8, i32 40, align 8 41 // CHECK: getelementptr {{.*}} i32 0 42 // CHECK: getelementptr {{.*}} i32 8 43 // CHECK: getelementptr {{.*}} i32 16 44 // CHECK: getelementptr {{.*}} i32 20 45 // CHECK: getelementptr {{.*}} i32 24 46 // CHECK: getelementptr {{.*}} i32 32 47 printf("%d %f %d %d %d %lld", 1, 2.0, 3, 4, 5, (long long)6); 48 } 49 50 __device__ void CheckNoArgs() { 51 // CHECK: call i32 @vprintf({{.*}}, i8* null){{$}} 52 printf("hello, world!"); 53 } 54