43 lines · plain
1// REQUIRES: x86-registered-target2// REQUIRES: nvptx-registered-target3 4// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -emit-llvm \5// RUN: -o - %s | FileCheck %s6 7#include "Inputs/cuda.h"8 9extern "C" __device__ int vprintf(const char*, const char*);10 11// Check a simple call to printf end-to-end.12// CHECK: [[SIMPLE_PRINTF_TY:%[a-zA-Z0-9_]+]] = type { i32, i64, double }13__device__ int CheckSimple() {14 // CHECK: [[BUF:%[a-zA-Z0-9_]+]] = alloca [[SIMPLE_PRINTF_TY]]15 // CHECK: [[FMT:%[0-9]+]] = load{{.*}}%fmt16 const char* fmt = "%d %lld %f";17 // CHECK: [[PTR0:%[0-9]+]] = getelementptr inbounds nuw [[SIMPLE_PRINTF_TY]], ptr [[BUF]], i32 0, i32 018 // CHECK: store i32 1, ptr [[PTR0]], align 419 // CHECK: [[PTR1:%[0-9]+]] = getelementptr inbounds nuw [[SIMPLE_PRINTF_TY]], ptr [[BUF]], i32 0, i32 120 // CHECK: store i64 2, ptr [[PTR1]], align 821 // CHECK: [[PTR2:%[0-9]+]] = getelementptr inbounds nuw [[SIMPLE_PRINTF_TY]], ptr [[BUF]], i32 0, i32 222 // CHECK: store double 3.0{{[^,]*}}, ptr [[PTR2]], align 823 // CHECK: [[RET:%[0-9]+]] = call i32 @vprintf(ptr [[FMT]], ptr [[BUF]])24 // CHECK: ret i32 [[RET]]25 return printf(fmt, 1, 2ll, 3.0);26}27 28__device__ void CheckNoArgs() {29 // CHECK: call i32 @vprintf({{.*}}, ptr null){{$}}30 printf("hello, world!");31}32 33// Check that printf's alloca happens in the entry block, not inside the if34// statement.35__device__ bool foo();36__device__ void CheckAllocaIsInEntryBlock() {37 // CHECK: alloca %printf_args38 // CHECK: call {{.*}} @_Z3foov()39 if (foo()) {40 printf("%d", 42);41 }42}43