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