1 // Test the Sema analysis of caller-callee relationships of host device 2 // functions when compiling CUDA code. There are 4 permutations of this test as 3 // host and device compilation are separate compilation passes, and clang has 4 // an option to allow host calls from host device functions. __CUDA_ARCH__ is 5 // defined when compiling for the device and TEST_WARN_HD when host calls are 6 // allowed from host device functions. So for example, if __CUDA_ARCH__ is 7 // defined and TEST_WARN_HD is not then device compilation is happening but 8 // host device functions are not allowed to call device functions. 9 10 // RUN: %clang_cc1 -fsyntax-only -verify %s 11 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -verify %s 12 // RUN: %clang_cc1 -fsyntax-only -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD 13 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -triple nvptx-unknown-cuda -fcuda-allow-host-calls-from-host-device -verify %s -DTEST_WARN_HD 14 15 #include "Inputs/cuda.h" 16 17 __host__ void hd1h(void); 18 #if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD) 19 // expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 20 #endif 21 __device__ void hd1d(void); 22 #ifndef __CUDA_ARCH__ 23 // expected-note@-2 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 24 #endif 25 __host__ void hd1hg(void); 26 __device__ void hd1dg(void); 27 #ifdef __CUDA_ARCH__ 28 __host__ void hd1hig(void); 29 #if !defined(TEST_WARN_HD) 30 // expected-note@-2 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 31 #endif 32 #else 33 __device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 34 #endif 35 __host__ __device__ void hd1hd(void); 36 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}} 37 38 __host__ __device__ void hd1(void) { 39 #if defined(TEST_WARN_HD) && defined(__CUDA_ARCH__) 40 // expected-warning@-2 {{calling __host__ function hd1h from __host__ __device__ function hd1}} 41 // expected-warning@-3 {{calling __host__ function hd1hig from __host__ __device__ function hd1}} 42 #endif 43 hd1d(); 44 #ifndef __CUDA_ARCH__ 45 // expected-error@-2 {{no matching function}} 46 #endif 47 hd1h(); 48 #if defined(__CUDA_ARCH__) && !defined(TEST_WARN_HD) 49 // expected-error@-2 {{no matching function}} 50 #endif 51 52 // No errors as guarded 53 #ifdef __CUDA_ARCH__ 54 hd1d(); 55 #else 56 hd1h(); 57 #endif 58 59 // Errors as incorrectly guarded 60 #ifndef __CUDA_ARCH__ 61 hd1dig(); // expected-error {{no matching function}} 62 #else 63 hd1hig(); 64 #ifndef TEST_WARN_HD 65 // expected-error@-2 {{no matching function}} 66 #endif 67 #endif 68 69 hd1hd(); 70 hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}} 71 } 72
