1 // RUN: %clang_cc1 -fsyntax-only -verify %s 2 3 // Check that it's OK for kernels to call HD functions that call device-only 4 // functions. 5 6 #include "Inputs/cuda.h" 7 device_fn(int)8__device__ void device_fn(int) {} 9 // expected-note@-1 2 {{declared here}} 10 hd1()11inline __host__ __device__ int hd1() { 12 device_fn(0); // expected-error {{reference to __device__ function}} 13 return 0; 14 } 15 hd2()16inline __host__ __device__ int hd2() { 17 // No error here because hd2 is only referenced from a kernel. 18 device_fn(0); 19 return 0; 20 } 21 hd3(int)22inline __host__ __device__ void hd3(int) { 23 device_fn(0); // expected-error {{reference to __device__ function 'device_fn'}} 24 } hd3(double)25inline __host__ __device__ void hd3(double) {} 26 hd4(int)27inline __host__ __device__ void hd4(int) {} hd4(double)28inline __host__ __device__ void hd4(double) { 29 device_fn(0); // No error; this function is never called. 30 } 31 kernel(int)32__global__ void kernel(int) { hd2(); } 33 34 template <typename T> launch_kernel()35void launch_kernel() { 36 kernel<<<0, 0>>>(T()); 37 38 // Notice that these two diagnostics are different: Because the call to hd1 39 // is not dependent on T, the call to hd1 comes from 'launch_kernel', while 40 // the call to hd3, being dependent, comes from 'launch_kernel<int>'. 41 hd1(); // expected-note {{called by 'launch_kernel<int>'}} 42 hd3(T()); // expected-note {{called by 'launch_kernel<int>'}} 43 } 44 host_fn()45void host_fn() { 46 launch_kernel<int>(); 47 // expected-note@-1 2 {{called by 'host_fn'}} 48 } 49