[CUDA] Report an error if code tries to mix incompatible CUDA attributes.

Summary: Thanks to jhen for helping me figure this out.

Reviewers: tra, echristo

Subscribers: jhen

Differential Revision: http://reviews.llvm.org/D16129

llvm-svn: 257554

3 files changed, 55 insertions, 1 deletions

diff --git a/clang/test/SemaCUDA/Inputs/cuda.h b/clang/test/SemaCUDA/Inputs/cuda.h
index a9a4595a14a..18cafdf96af 100644
--- a/clang/test/SemaCUDA/Inputs/cuda.h
+++ b/clang/test/SemaCUDA/Inputs/cuda.h
@@ -2,6 +2,9 @@
 
 #include <stddef.h>
 
+// Make this file work with nvcc, for testing compatibility.
+
+#ifndef __NVCC__
 #define __constant__ __attribute__((constant))
 #define __device__ __attribute__((device))
 #define __global__ __attribute__((global))
@@ -18,3 +21,4 @@ typedef struct cudaStream *cudaStream_t;
 
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
+#endif // !__NVCC__
diff --git a/clang/test/SemaCUDA/attributes.cu b/clang/test/SemaCUDA/attributes-on-non-cuda.cu
index ce4dc925a3f..e9e32ce658c 100644
--- a/clang/test/SemaCUDA/attributes.cu
+++ b/clang/test/SemaCUDA/attributes-on-non-cuda.cu
@@ -1,4 +1,5 @@
-// Tests handling of CUDA attributes.
+// Tests that CUDA attributes are warnings when compiling C files, but not when
+// compiling CUDA files.
 //
 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s
diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
new file mode 100644
index 00000000000..7e01e141de1
--- /dev/null
+++ b/clang/test/SemaCUDA/bad-attributes.cu
@@ -0,0 +1,49 @@
+// Tests handling of CUDA attributes that are bad either because they're
+// applied to the wrong sort of thing, or because they're given in illegal
+// combinations.
+//
+// You should be able to run this file through nvcc for compatibility testing.
+//
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+#include "Inputs/cuda.h"
+
+// Try applying attributes to functions and variables.  Some should generate
+// warnings; others not.
+__device__ int a1;
+__device__ void a2();
+__host__ int b1; // expected-warning {{attribute only applies to functions}}
+__host__ void b2();
+__constant__ int c1;
+__constant__ void c2(); // expected-warning {{attribute only applies to variables}}
+__shared__ int d1;
+__shared__ void d2(); // expected-warning {{attribute only applies to variables}}
+__global__ int e1; // expected-warning {{attribute only applies to functions}}
+__global__ void e2();
+
+// Try all pairs of attributes which can be present on a function or a
+// variable.  Check both orderings of the attributes, as that can matter in
+// clang.
+__device__ __host__ void z1();
+__device__ __constant__ int z2;
+__device__ __shared__ int z3;
+__device__ __global__ void z4(); // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+
+__host__ __device__ void z5();
+__host__ __global__ void z6();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+
+__constant__ __device__ int z7;
+__constant__ __shared__ int z8;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+
+__shared__ __device__ int z9;
+__shared__ __constant__ int z10;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+
+__global__ __device__ void z11();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__global__ __host__ void z12();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}