r257554 - [CUDA] Report an error if code tries to mix incompatible CUDA attributes.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Tue Jan 12 17:07:35 PST 2016
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 at +12 {{'device' attribute ignored}}
-// expected-warning at +12 {{'global' attribute ignored}}
-// expected-warning at +12 {{'constant' attribute ignored}}
-// expected-warning at +12 {{'shared' attribute ignored}}
-// expected-warning at +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 at -1 {{conflicting attribute is here}}
+
+__host__ __device__ void z5();
+__host__ __global__ void z6(); // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
+
+__constant__ __device__ int z7;
+__constant__ __shared__ int z8; // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
+
+__shared__ __device__ int z9;
+__shared__ __constant__ int z10; // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
+
+__global__ __device__ void z11(); // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
+__global__ __host__ void z12(); // expected-error {{attributes are not compatible}}
+// expected-note at -1 {{conflicting attribute is here}}
More information about the cfe-commits
mailing list