1*67e74705SXin Li // Tests handling of CUDA attributes that are bad either because they're 2*67e74705SXin Li // applied to the wrong sort of thing, or because they're given in illegal 3*67e74705SXin Li // combinations. 4*67e74705SXin Li // 5*67e74705SXin Li // You should be able to run this file through nvcc for compatibility testing. 6*67e74705SXin Li // 7*67e74705SXin Li // RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s 8*67e74705SXin Li // RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s 9*67e74705SXin Li 10*67e74705SXin Li #include "Inputs/cuda.h" 11*67e74705SXin Li 12*67e74705SXin Li // Try applying attributes to functions and variables. Some should generate 13*67e74705SXin Li // warnings; others not. 14*67e74705SXin Li __device__ int a1; 15*67e74705SXin Li __device__ void a2(); 16*67e74705SXin Li __host__ int b1; // expected-warning {{attribute only applies to functions}} 17*67e74705SXin Li __host__ void b2(); 18*67e74705SXin Li __constant__ int c1; 19*67e74705SXin Li __constant__ void c2(); // expected-warning {{attribute only applies to variables}} 20*67e74705SXin Li __shared__ int d1; 21*67e74705SXin Li __shared__ void d2(); // expected-warning {{attribute only applies to variables}} 22*67e74705SXin Li __global__ int e1; // expected-warning {{attribute only applies to functions}} 23*67e74705SXin Li __global__ void e2(); 24*67e74705SXin Li 25*67e74705SXin Li // Try all pairs of attributes which can be present on a function or a 26*67e74705SXin Li // variable. Check both orderings of the attributes, as that can matter in 27*67e74705SXin Li // clang. 28*67e74705SXin Li __device__ __host__ void z1(); 29*67e74705SXin Li __device__ __constant__ int z2; 30*67e74705SXin Li __device__ __shared__ int z3; 31*67e74705SXin Li __device__ __global__ void z4(); // expected-error {{attributes are not compatible}} 32*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 33*67e74705SXin Li 34*67e74705SXin Li __host__ __device__ void z5(); 35*67e74705SXin Li __host__ __global__ void z6(); // expected-error {{attributes are not compatible}} 36*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 37*67e74705SXin Li 38*67e74705SXin Li __constant__ __device__ int z7; 39*67e74705SXin Li __constant__ __shared__ int z8; // expected-error {{attributes are not compatible}} 40*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 41*67e74705SXin Li 42*67e74705SXin Li __shared__ __device__ int z9; 43*67e74705SXin Li __shared__ __constant__ int z10; // expected-error {{attributes are not compatible}} 44*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 45*67e74705SXin Li 46*67e74705SXin Li __global__ __device__ void z11(); // expected-error {{attributes are not compatible}} 47*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 48*67e74705SXin Li __global__ __host__ void z12(); // expected-error {{attributes are not compatible}} 49*67e74705SXin Li // expected-note@-1 {{conflicting attribute is here}} 50*67e74705SXin Li 51*67e74705SXin Li struct S { fooS52*67e74705SXin Li __global__ void foo() {}; // expected-error {{must be a free function or static member function}} 53*67e74705SXin Li __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}} 54*67e74705SXin Li // Although this is implicitly inline, we shouldn't warn. bazS55*67e74705SXin Li __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}} 56*67e74705SXin Li }; 57*67e74705SXin Li foobar()58*67e74705SXin Li__global__ static inline void foobar() {}; 59*67e74705SXin Li #ifdef EXPECT_INLINE_WARNING 60*67e74705SXin Li // expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}} 61*67e74705SXin Li #endif 62