From a9b21a2f1a6122a82dfff6ffe9f21927d81fa22d Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Wed, 13 Jan 2016 01:07:35 +0000 Subject: [PATCH] [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 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@257554 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Sema/SemaDeclAttr.cpp | 48 +++++++++++++----- test/SemaCUDA/Inputs/cuda.h | 4 ++ ...ttributes.cu => attributes-on-non-cuda.cu} | 3 +- test/SemaCUDA/bad-attributes.cu | 49 +++++++++++++++++++ 4 files changed, 90 insertions(+), 14 deletions(-) rename test/SemaCUDA/{attributes.cu => attributes-on-non-cuda.cu} (93%) create mode 100644 test/SemaCUDA/bad-attributes.cu diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp index 5a0f0f84af..f94c822b90 100644 --- a/lib/Sema/SemaDeclAttr.cpp +++ b/lib/Sema/SemaDeclAttr.cpp @@ -348,6 +348,25 @@ static void handleSimpleAttribute(Sema &S, Decl *D, 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 handleOptimizeNoneAttr(Sema &S, Decl *D, } 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 @@ static void handleInterruptAttr(Sema &S, Decl *D, const AttributeList &Attr) { 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 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, 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 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, 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 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, 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 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, 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); diff --git a/test/SemaCUDA/Inputs/cuda.h b/test/SemaCUDA/Inputs/cuda.h index a9a4595a14..18cafdf96a 100644 --- a/test/SemaCUDA/Inputs/cuda.h +++ b/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 @@ typedef struct cudaStream *cudaStream_t; int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, cudaStream_t stream = 0); +#endif // !__NVCC__ diff --git a/test/SemaCUDA/attributes.cu b/test/SemaCUDA/attributes-on-non-cuda.cu similarity index 93% rename from test/SemaCUDA/attributes.cu rename to test/SemaCUDA/attributes-on-non-cuda.cu index ce4dc925a3..e9e32ce658 100644 --- a/test/SemaCUDA/attributes.cu +++ b/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/test/SemaCUDA/bad-attributes.cu b/test/SemaCUDA/bad-attributes.cu new file mode 100644 index 0000000000..7e01e141de --- /dev/null +++ b/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}} -- 2.40.0