1 // RUN: %clang_cc1 -fsyntax-only -verify %s 2 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s 3 4 #include "Inputs/cuda.h" 5 6 __host__ void h1h(void); 7 __device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}} 8 __host__ __device__ void h1hd(void); 9 __global__ void h1g(void); 10 11 struct h1ds { // expected-note {{requires 1 argument}} 12 __device__ h1ds(); // expected-note {{candidate constructor not viable: call to __device__ function from __host__ function}} 13 }; 14 h1(void)15__host__ void h1(void) { 16 h1h(); 17 h1d(); // expected-error {{no matching function}} 18 h1hd(); 19 h1g<<<1, 1>>>(); 20 h1ds x; // expected-error {{no matching constructor}} 21 } 22 23 __host__ void d1h(void); // expected-note {{candidate function not viable: call to __host__ function from __device__ function}} 24 __device__ void d1d(void); 25 __host__ __device__ void d1hd(void); 26 __global__ void d1g(void); // expected-note {{'d1g' declared here}} 27 d1(void)28__device__ void d1(void) { 29 d1h(); // expected-error {{no matching function}} 30 d1d(); 31 d1hd(); 32 d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}} 33 } 34 35 // Expected 0-1 as in one of host/device side compilation it is an error, while 36 // not in the other 37 __host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 38 __device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 39 __host__ void hd1hg(void); 40 __device__ void hd1dg(void); 41 #ifdef __CUDA_ARCH__ 42 __host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}} 43 #else 44 __device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}} 45 #endif 46 __host__ __device__ void hd1hd(void); 47 __global__ void hd1g(void); // expected-note {{'hd1g' declared here}} 48 hd1(void)49__host__ __device__ void hd1(void) { 50 // Expected 0-1 as in one of host/device side compilation it is an error, 51 // while not in the other 52 hd1d(); // expected-error 0-1 {{no matching function}} 53 hd1h(); // expected-error 0-1 {{no matching function}} 54 55 // No errors as guarded 56 #ifdef __CUDA_ARCH__ 57 hd1d(); 58 #else 59 hd1h(); 60 #endif 61 62 // Errors as incorrectly guarded 63 #ifndef __CUDA_ARCH__ 64 hd1dig(); // expected-error {{no matching function}} 65 #else 66 hd1hig(); // expected-error {{no matching function}} 67 #endif 68 69 hd1hd(); 70 hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}} 71 } 72