blob: 8140f619361bc889dc2d8b8c6b28bf936bac1dfd [file] [log] [blame]
Stephen Hines0e2c34f2015-03-23 12:09:02 -07001// 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
2
3#include "Inputs/cuda.h"
4
5extern "C"
6void host_function() {}
7
8// CHECK-LABEL: define void @hd_function_a
9extern "C"
10__host__ __device__ void hd_function_a() {
11 // CHECK: call void @host_function
12 host_function();
13}
14
15// CHECK: declare void @host_function
16
17// CHECK-LABEL: define void @hd_function_b
18extern "C"
19__host__ __device__ void hd_function_b(bool b) { if (b) host_function(); }
20
21// CHECK-LABEL: define void @device_function_b
22extern "C"
23__device__ void device_function_b() { hd_function_b(false); }
24
25// CHECK-LABEL: define void @global_function
26extern "C"
27__global__ void global_function() {
28 // CHECK: call void @device_function_b
29 device_function_b();
30}
31
32// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}