diff options
Diffstat (limited to 'test/CodeGenCUDA/host-device-calls-host.cu')
-rw-r--r-- | test/CodeGenCUDA/host-device-calls-host.cu | 32 |
1 files changed, 32 insertions, 0 deletions
diff --git a/test/CodeGenCUDA/host-device-calls-host.cu b/test/CodeGenCUDA/host-device-calls-host.cu new file mode 100644 index 0000000..8140f61 --- /dev/null +++ b/test/CodeGenCUDA/host-device-calls-host.cu @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-allow-host-calls-from-host-device -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s + +#include "Inputs/cuda.h" + +extern "C" +void host_function() {} + +// CHECK-LABEL: define void @hd_function_a +extern "C" +__host__ __device__ void hd_function_a() { + // CHECK: call void @host_function + host_function(); +} + +// CHECK: declare void @host_function + +// CHECK-LABEL: define void @hd_function_b +extern "C" +__host__ __device__ void hd_function_b(bool b) { if (b) host_function(); } + +// CHECK-LABEL: define void @device_function_b +extern "C" +__device__ void device_function_b() { hd_function_b(false); } + +// CHECK-LABEL: define void @global_function +extern "C" +__global__ void global_function() { + // CHECK: call void @device_function_b + device_function_b(); +} + +// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1} |