| 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 | |
|---|
| 8 | __device__ void device_fn(int) {} |
|---|
| 9 | // expected-note@-1 2 {{declared here}} |
|---|
| 10 | |
|---|
| 11 | inline __host__ __device__ int hd1() { |
|---|
| 12 | device_fn(0); // expected-error {{reference to __device__ function}} |
|---|
| 13 | return 0; |
|---|
| 14 | } |
|---|
| 15 | |
|---|
| 16 | inline __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 | |
|---|
| 22 | inline __host__ __device__ void hd3(int) { |
|---|
| 23 | device_fn(0); // expected-error {{reference to __device__ function 'device_fn'}} |
|---|
| 24 | } |
|---|
| 25 | inline __host__ __device__ void hd3(double) {} |
|---|
| 26 | |
|---|
| 27 | inline __host__ __device__ void hd4(int) {} |
|---|
| 28 | inline __host__ __device__ void hd4(double) { |
|---|
| 29 | device_fn(0); // No error; this function is never called. |
|---|
| 30 | } |
|---|
| 31 | |
|---|
| 32 | __global__ void kernel(int) { hd2(); } |
|---|
| 33 | |
|---|
| 34 | template <typename T> |
|---|
| 35 | void 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'}} |
|---|
| 42 | hd3(T()); // expected-note {{called by 'launch_kernel<int>'}} |
|---|
| 43 | } |
|---|
| 44 | |
|---|
| 45 | void host_fn() { |
|---|
| 46 | launch_kernel<int>(); |
|---|
| 47 | // expected-note@-1 2 {{called by 'host_fn'}} |
|---|
| 48 | } |
|---|
| 49 | |
|---|