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