diff options
Diffstat (limited to 'test/SemaCUDA/function-target.cu')
-rw-r--r-- | test/SemaCUDA/function-target.cu | 35 |
1 files changed, 31 insertions, 4 deletions
diff --git a/test/SemaCUDA/function-target.cu b/test/SemaCUDA/function-target.cu index 51bc8c9..ca56030 100644 --- a/test/SemaCUDA/function-target.cu +++ b/test/SemaCUDA/function-target.cu @@ -1,4 +1,5 @@ // RUN: %clang_cc1 -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s #include "Inputs/cuda.h" @@ -31,14 +32,40 @@ __device__ void d1(void) { d1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'd1g' in __device__ function}} } -__host__ void hd1h(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}} -__device__ void hd1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}} +// Expected 0-1 as in one of host/device side compilation it is an error, while +// not in the other +__host__ void hd1h(void); // expected-note 0-1 {{candidate function not viable: call to __host__ function from __host__ __device__ function}} +__device__ void hd1d(void); // expected-note 0-1 {{candidate function not viable: call to __device__ function from __host__ __device__ function}} +__host__ void hd1hg(void); +__device__ void hd1dg(void); +#ifdef __CUDA_ARCH__ +__host__ void hd1hig(void); // expected-note {{candidate function not viable: call to __host__ function from __host__ __device__ function}} +#else +__device__ void hd1dig(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ __device__ function}} +#endif __host__ __device__ void hd1hd(void); __global__ void hd1g(void); // expected-note {{'hd1g' declared here}} __host__ __device__ void hd1(void) { - hd1h(); // expected-error {{no matching function}} - hd1d(); // expected-error {{no matching function}} + // Expected 0-1 as in one of host/device side compilation it is an error, + // while not in the other + hd1d(); // expected-error 0-1 {{no matching function}} + hd1h(); // expected-error 0-1 {{no matching function}} + + // No errors as guarded +#ifdef __CUDA_ARCH__ + hd1d(); +#else + hd1h(); +#endif + + // Errors as incorrectly guarded +#ifndef __CUDA_ARCH__ + hd1dig(); // expected-error {{no matching function}} +#else + hd1hig(); // expected-error {{no matching function}} +#endif + hd1hd(); hd1g<<<1, 1>>>(); // expected-error {{reference to __global__ function 'hd1g' in __host__ __device__ function}} } |