Index: cfe/trunk/lib/Sema/SemaDeclAttr.cpp =================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp @@ -348,6 +348,25 @@ Attr.getAttributeSpellingListIndex())); } +template +static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D, + const AttributeList &Attr) { + handleSimpleAttribute(S, D, Attr); +} + +/// \brief Applies the given attribute to the Decl so long as the Decl doesn't +/// already have one of the given incompatible attributes. +template +static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D, + const AttributeList &Attr) { + if (checkAttrMutualExclusion(S, D, Attr.getRange(), + Attr.getName())) + return; + handleSimpleAttributeWithExclusions(S, D, + Attr); +} + /// \brief Check if the passed-in expression is of type int or bool. static bool isIntOrBool(Expr *Exp) { QualType QT = Exp->getType(); @@ -3588,6 +3607,12 @@ } static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (checkAttrMutualExclusion(S, D, Attr.getRange(), + Attr.getName()) || + checkAttrMutualExclusion(S, D, Attr.getRange(), + Attr.getName())) { + return; + } FunctionDecl *FD = cast(D); if (!FD->getReturnType()->isVoidType()) { SourceRange RTRange = FD->getReturnTypeSourceRange(); @@ -4558,14 +4583,6 @@ handleARMInterruptAttr(S, D, Attr); } -static void handleMips16Attribute(Sema &S, Decl *D, const AttributeList &Attr) { - if (checkAttrMutualExclusion(S, D, Attr.getRange(), - Attr.getName())) - return; - - handleSimpleAttribute(S, D, Attr); -} - static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D, const AttributeList &Attr) { uint32_t NumRegs; @@ -4955,7 +4972,8 @@ handleDLLAttr(S, D, Attr); break; case AttributeList::AT_Mips16: - handleMips16Attribute(S, D, Attr); + handleSimpleAttributeWithExclusions(S, D, + Attr); break; case AttributeList::AT_NoMips16: handleSimpleAttribute(S, D, Attr); @@ -5006,7 +5024,8 @@ handleCommonAttr(S, D, Attr); break; case AttributeList::AT_CUDAConstant: - handleSimpleAttribute(S, D, Attr); + handleSimpleAttributeWithExclusions(S, D, + Attr); break; case AttributeList::AT_PassObjectSize: handlePassObjectSizeAttr(S, D, Attr); @@ -5051,10 +5070,12 @@ handleGlobalAttr(S, D, Attr); break; case AttributeList::AT_CUDADevice: - handleSimpleAttribute(S, D, Attr); + handleSimpleAttributeWithExclusions(S, D, + Attr); break; case AttributeList::AT_CUDAHost: - handleSimpleAttribute(S, D, Attr); + handleSimpleAttributeWithExclusions(S, D, + Attr); break; case AttributeList::AT_GNUInline: handleGNUInlineAttr(S, D, Attr); @@ -5114,7 +5135,8 @@ handleSimpleAttribute(S, D, Attr); break; case AttributeList::AT_CUDAShared: - handleSimpleAttribute(S, D, Attr); + handleSimpleAttributeWithExclusions(S, D, + Attr); break; case AttributeList::AT_VecReturn: handleVecReturnAttr(S, D, Attr); Index: cfe/trunk/test/SemaCUDA/Inputs/cuda.h =================================================================== --- cfe/trunk/test/SemaCUDA/Inputs/cuda.h +++ cfe/trunk/test/SemaCUDA/Inputs/cuda.h @@ -2,6 +2,9 @@ #include +// 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 @@ int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, cudaStream_t stream = 0); +#endif // !__NVCC__ Index: cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu =================================================================== --- cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu +++ cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu @@ -0,0 +1,34 @@ +// 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 +// Now pretend that we're compiling a C file. There should be warnings. +// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s + +#if defined(EXPECT_WARNINGS) +// expected-warning@+12 {{'device' attribute ignored}} +// expected-warning@+12 {{'global' attribute ignored}} +// expected-warning@+12 {{'constant' attribute ignored}} +// expected-warning@+12 {{'shared' attribute ignored}} +// expected-warning@+12 {{'host' attribute ignored}} +// +// NOTE: IgnoredAttr in clang which is used for the rest of +// attributes ignores LangOpts, so there are no warnings. +#else +// expected-no-diagnostics +#endif + +__attribute__((device)) void f_device(); +__attribute__((global)) void f_global(); +__attribute__((constant)) int* g_constant; +__attribute__((shared)) float *g_shared; +__attribute__((host)) void f_host(); +__attribute__((device_builtin)) void f_device_builtin(); +typedef __attribute__((device_builtin)) const void *t_device_builtin; +enum __attribute__((device_builtin)) e_device_builtin {E}; +__attribute__((device_builtin)) int v_device_builtin; +__attribute__((cudart_builtin)) void f_cudart_builtin(); +__attribute__((nv_weak)) void f_nv_weak(); +__attribute__((device_builtin_surface_type)) unsigned long long surface_var; +__attribute__((device_builtin_texture_type)) unsigned long long texture_var; Index: cfe/trunk/test/SemaCUDA/attributes.cu =================================================================== --- cfe/trunk/test/SemaCUDA/attributes.cu +++ cfe/trunk/test/SemaCUDA/attributes.cu @@ -1,33 +0,0 @@ -// Tests handling of CUDA attributes. -// -// RUN: %clang_cc1 -fsyntax-only -verify %s -// RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify %s -// Now pretend that we're compiling a C file. There should be warnings. -// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s - -#if defined(EXPECT_WARNINGS) -// expected-warning@+12 {{'device' attribute ignored}} -// expected-warning@+12 {{'global' attribute ignored}} -// expected-warning@+12 {{'constant' attribute ignored}} -// expected-warning@+12 {{'shared' attribute ignored}} -// expected-warning@+12 {{'host' attribute ignored}} -// -// NOTE: IgnoredAttr in clang which is used for the rest of -// attributes ignores LangOpts, so there are no warnings. -#else -// expected-no-diagnostics -#endif - -__attribute__((device)) void f_device(); -__attribute__((global)) void f_global(); -__attribute__((constant)) int* g_constant; -__attribute__((shared)) float *g_shared; -__attribute__((host)) void f_host(); -__attribute__((device_builtin)) void f_device_builtin(); -typedef __attribute__((device_builtin)) const void *t_device_builtin; -enum __attribute__((device_builtin)) e_device_builtin {E}; -__attribute__((device_builtin)) int v_device_builtin; -__attribute__((cudart_builtin)) void f_cudart_builtin(); -__attribute__((nv_weak)) void f_nv_weak(); -__attribute__((device_builtin_surface_type)) unsigned long long surface_var; -__attribute__((device_builtin_texture_type)) unsigned long long texture_var; Index: cfe/trunk/test/SemaCUDA/bad-attributes.cu =================================================================== --- cfe/trunk/test/SemaCUDA/bad-attributes.cu +++ cfe/trunk/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}}