| 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 | __constant__ __shared__ int z10a; // expected-error {{attributes are not compatible}} |
|---|
| 46 | // expected-note@-1 {{conflicting attribute is here}} |
|---|
| 47 | |
|---|
| 48 | __global__ __device__ void z11(); // expected-error {{attributes are not compatible}} |
|---|
| 49 | // expected-note@-1 {{conflicting attribute is here}} |
|---|
| 50 | __global__ __host__ void z12(); // expected-error {{attributes are not compatible}} |
|---|
| 51 | // expected-note@-1 {{conflicting attribute is here}} |
|---|
| 52 | |
|---|
| 53 | struct S { |
|---|
| 54 | __global__ void foo() {}; // expected-error {{must be a free function or static member function}} |
|---|
| 55 | __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}} |
|---|
| 56 | // Although this is implicitly inline, we shouldn't warn. |
|---|
| 57 | __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}} |
|---|
| 58 | }; |
|---|
| 59 | |
|---|
| 60 | __global__ static inline void foobar() {}; |
|---|
| 61 | #ifdef EXPECT_INLINE_WARNING |
|---|
| 62 | // expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}} |
|---|
| 63 | #endif |
|---|
| 64 | |
|---|
| 65 | __constant__ int global_constant; |
|---|
| 66 | void host_fn() { |
|---|
| 67 | __constant__ int c; // expected-error {{__constant__ variables must be global}} |
|---|
| 68 | __shared__ int s; // expected-error {{__shared__ local variables not allowed in __host__ functions}} |
|---|
| 69 | } |
|---|
| 70 | __device__ void device_fn() { |
|---|
| 71 | __constant__ int c; // expected-error {{__constant__ variables must be global}} |
|---|
| 72 | } |
|---|
| 73 | |
|---|