Author: jlebar Date: Tue Jan 12 19:07:35 2016 New Revision: 257554 URL: http://llvm.org/viewvc/llvm-project?rev=257554&view=rev Log: [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 Added: cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu - copied, changed from r257543, cfe/trunk/test/SemaCUDA/attributes.cu cfe/trunk/test/SemaCUDA/bad-attributes.cu Removed: cfe/trunk/test/SemaCUDA/attributes.cu Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp cfe/trunk/test/SemaCUDA/Inputs/cuda.h Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=257554&r1=257553&r2=257554&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original) +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Tue Jan 12 19:07:35 2016 @@ -348,6 +348,25 @@ static void handleSimpleAttribute(Sema & Attr.getAttributeSpellingListIndex())); } +template <typename AttrType> +static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D, + const AttributeList &Attr) { + handleSimpleAttribute<AttrType>(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 <typename AttrType, typename IncompatibleAttrType, + typename... IncompatibleAttrTypes> +static void handleSimpleAttributeWithExclusions(Sema &S, Decl *D, + const AttributeList &Attr) { + if (checkAttrMutualExclusion<IncompatibleAttrType>(S, D, Attr.getRange(), + Attr.getName())) + return; + handleSimpleAttributeWithExclusions<AttrType, IncompatibleAttrTypes...>(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 handleOptimizeNoneAttr(Sema } static void handleGlobalAttr(Sema &S, Decl *D, const AttributeList &Attr) { + if (checkAttrMutualExclusion<CUDADeviceAttr>(S, D, Attr.getRange(), + Attr.getName()) || + checkAttrMutualExclusion<CUDAHostAttr>(S, D, Attr.getRange(), + Attr.getName())) { + return; + } FunctionDecl *FD = cast<FunctionDecl>(D); if (!FD->getReturnType()->isVoidType()) { SourceRange RTRange = FD->getReturnTypeSourceRange(); @@ -4558,14 +4583,6 @@ static void handleInterruptAttr(Sema &S, handleARMInterruptAttr(S, D, Attr); } -static void handleMips16Attribute(Sema &S, Decl *D, const AttributeList &Attr) { - if (checkAttrMutualExclusion<MipsInterruptAttr>(S, D, Attr.getRange(), - Attr.getName())) - return; - - handleSimpleAttribute<Mips16Attr>(S, D, Attr); -} - static void handleAMDGPUNumVGPRAttr(Sema &S, Decl *D, const AttributeList &Attr) { uint32_t NumRegs; @@ -4955,7 +4972,8 @@ static void ProcessDeclAttribute(Sema &S handleDLLAttr(S, D, Attr); break; case AttributeList::AT_Mips16: - handleMips16Attribute(S, D, Attr); + handleSimpleAttributeWithExclusions<Mips16Attr, MipsInterruptAttr>(S, D, + Attr); break; case AttributeList::AT_NoMips16: handleSimpleAttribute<NoMips16Attr>(S, D, Attr); @@ -5006,7 +5024,8 @@ static void ProcessDeclAttribute(Sema &S handleCommonAttr(S, D, Attr); break; case AttributeList::AT_CUDAConstant: - handleSimpleAttribute<CUDAConstantAttr>(S, D, Attr); + handleSimpleAttributeWithExclusions<CUDAConstantAttr, CUDASharedAttr>(S, D, + Attr); break; case AttributeList::AT_PassObjectSize: handlePassObjectSizeAttr(S, D, Attr); @@ -5051,10 +5070,12 @@ static void ProcessDeclAttribute(Sema &S handleGlobalAttr(S, D, Attr); break; case AttributeList::AT_CUDADevice: - handleSimpleAttribute<CUDADeviceAttr>(S, D, Attr); + handleSimpleAttributeWithExclusions<CUDADeviceAttr, CUDAGlobalAttr>(S, D, + Attr); break; case AttributeList::AT_CUDAHost: - handleSimpleAttribute<CUDAHostAttr>(S, D, Attr); + handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, + Attr); break; case AttributeList::AT_GNUInline: handleGNUInlineAttr(S, D, Attr); @@ -5114,7 +5135,8 @@ static void ProcessDeclAttribute(Sema &S handleSimpleAttribute<NoThrowAttr>(S, D, Attr); break; case AttributeList::AT_CUDAShared: - handleSimpleAttribute<CUDASharedAttr>(S, D, Attr); + handleSimpleAttributeWithExclusions<CUDASharedAttr, CUDAConstantAttr>(S, D, + Attr); break; case AttributeList::AT_VecReturn: handleVecReturnAttr(S, D, Attr); Modified: cfe/trunk/test/SemaCUDA/Inputs/cuda.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/Inputs/cuda.h?rev=257554&r1=257553&r2=257554&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/Inputs/cuda.h (original) +++ cfe/trunk/test/SemaCUDA/Inputs/cuda.h Tue Jan 12 19:07:35 2016 @@ -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__ Copied: cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu (from r257543, cfe/trunk/test/SemaCUDA/attributes.cu) URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu?p2=cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu&p1=cfe/trunk/test/SemaCUDA/attributes.cu&r1=257543&r2=257554&rev=257554&view=diff ============================================================================== --- cfe/trunk/test/SemaCUDA/attributes.cu (original) +++ cfe/trunk/test/SemaCUDA/attributes-on-non-cuda.cu Tue Jan 12 19:07:35 2016 @@ -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 Removed: cfe/trunk/test/SemaCUDA/attributes.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/attributes.cu?rev=257553&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/attributes.cu (original) +++ cfe/trunk/test/SemaCUDA/attributes.cu (removed) @@ -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; Added: cfe/trunk/test/SemaCUDA/bad-attributes.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/bad-attributes.cu?rev=257554&view=auto ============================================================================== --- cfe/trunk/test/SemaCUDA/bad-attributes.cu (added) +++ cfe/trunk/test/SemaCUDA/bad-attributes.cu Tue Jan 12 19:07:35 2016 @@ -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}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits