xref: /aosp_15_r20/external/clang/test/CodeGenCUDA/host-device-calls-host.cu (revision 67e74705e28f6214e480b399dd47ea732279e315) 
1*67e74705SXin Li // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
2*67e74705SXin Li 
3*67e74705SXin Li #include "Inputs/cuda.h"
4*67e74705SXin Li 
5*67e74705SXin Li extern "C"
host_function()6*67e74705SXin Li void host_function() {}
7*67e74705SXin Li 
8*67e74705SXin Li // CHECK-LABEL: define void @hd_function_a
9*67e74705SXin Li extern "C"
hd_function_a()10*67e74705SXin Li __host__ __device__ void hd_function_a() {
11*67e74705SXin Li   // CHECK: call void @host_function
12*67e74705SXin Li   host_function();
13*67e74705SXin Li }
14*67e74705SXin Li 
15*67e74705SXin Li // CHECK: declare void @host_function
16*67e74705SXin Li 
17*67e74705SXin Li // CHECK-LABEL: define void @hd_function_b
18*67e74705SXin Li extern "C"
hd_function_b(bool b)19*67e74705SXin Li __host__ __device__ void hd_function_b(bool b) { if (b) host_function(); }
20*67e74705SXin Li 
21*67e74705SXin Li // CHECK-LABEL: define void @device_function_b
22*67e74705SXin Li extern "C"
device_function_b()23*67e74705SXin Li __device__ void device_function_b() { hd_function_b(false); }
24*67e74705SXin Li 
25*67e74705SXin Li // CHECK-LABEL: define void @global_function
26*67e74705SXin Li extern "C"
global_function()27*67e74705SXin Li __global__ void global_function() {
28*67e74705SXin Li   // CHECK: call void @device_function_b
29*67e74705SXin Li   device_function_b();
30*67e74705SXin Li }
31*67e74705SXin Li 
32*67e74705SXin Li // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
33