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

[jlebar]

Summary

Thanks to jhen for helping me figure this out.

[tra]

lib/Sema/SemaDeclAttr.cpp

364

No need for {}

test/SemaCUDA/bad-attributes.cu

12–20

Those are already defined in Inputs/cuda.h in the test directory:

#include "Inputs/cuda.h"

34–35

Those are unsupported stubs that exist only to avoid warnings when we include CUDA headers.
No point testing for them here IMO.

[jlebar]

Address tra's review comments.

[jlebar]

Thank you for the review.

lib/Sema/SemaDeclAttr.cpp

364

For once this was intentional; this if body is multiline. See also handleGlobalAttr below.

I'm happy to do whatever the style is, but I really hope it doesn't prescribe no braces here.

test/SemaCUDA/bad-attributes.cu

34–35

Done. I presume the same applies to cudart_builtin, nv_weak, device_builtin_surface_type, and device_builtin_texture_type?

[tra]

LGTM

lib/Sema/SemaDeclAttr.cpp

 /// additional semantic checking.
 template <typename AttrType>
 static void handleSimpleAttribute(Sema &S, Decl *D,
                                   const AttributeList &Attr) {
   D->addAttr(::new (S.Context) AttrType(Attr.getRange(), S.Context,
                                         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();
   return QT->isBooleanType() || QT->isIntegerType();
 }
 
 
 // Check to see if the type is a smart pointer of some kind.  We assume

 static void handleOptimizeNoneAttr(Sema &S, Decl *D,
                                    const AttributeList &Attr) {
   if (OptimizeNoneAttr *Optnone = S.mergeOptimizeNoneAttr(
           D, Attr.getRange(), Attr.getAttributeSpellingListIndex()))
     D->addAttr(Optnone);
 }
 
 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();
     S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
         << FD->getType()
         << (RTRange.isValid() ? FixItHint::CreateReplacement(RTRange, "void")
                               : FixItHint());
     return;
else if (S.Context.getTargetInfo().getTriple().getArch() ==
                llvm::Triple::mipsel ||
            S.Context.getTargetInfo().getTriple().getArch() ==
                llvm::Triple::mips)
     handleMipsInterruptAttr(S, D, Attr);
   else
     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;
   Expr *NumRegsExpr = static_cast<Expr *>(Attr.getArgAsExpr(0));
   if (!checkUInt32Argument(S, Attr, NumRegsExpr, NumRegs))
     return;
 
   D->addAttr(::new (S.Context)
static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case AttributeList::AT_X86ForceAlignArgPointer:
     handleX86ForceAlignArgPointerAttr(S, D, Attr);
     break;
   case AttributeList::AT_DLLExport:
   case AttributeList::AT_DLLImport:
     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);
     break;
   case AttributeList::AT_AMDGPUNumVGPR:
     handleAMDGPUNumVGPRAttr(S, D, Attr);
     break;
   case AttributeList::AT_AMDGPUNumSGPR:
case AttributeList::AT_Availability:
     break;
   case AttributeList::AT_CarriesDependency:
     handleDependencyAttr(S, scope, D, Attr);
     break;
   case AttributeList::AT_Common:
     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);
     break;
   case AttributeList::AT_Constructor:
     handleConstructorAttr(S, D, Attr);
     break;
   case AttributeList::AT_CXX11NoReturn:
case AttributeList::AT_Format:
     break;
   case AttributeList::AT_FormatArg:
     handleFormatArgAttr(S, D, Attr);
     break;
   case AttributeList::AT_CUDAGlobal:
     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);
     break;
   case AttributeList::AT_CUDALaunchBounds:
     handleLaunchBoundsAttr(S, D, Attr);
     break;
   case AttributeList::AT_Restrict:
case AttributeList::AT_Naked:
     break;
   case AttributeList::AT_NoReturn:
     handleNoReturnAttr(S, D, Attr);
     break;
   case AttributeList::AT_NoThrow:
     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);
     break;
 
   case AttributeList::AT_ObjCOwnership:
     handleObjCOwnershipAttr(S, D, Attr);
     break;

test/SemaCUDA/Inputs/cuda.h

 /* Minimal declarations for CUDA support.  Testing purposes only. */
 
 #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))
 #define __host__ __attribute__((host))
 #define __shared__ __attribute__((shared))
 #define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__)))
 
 struct dim3 {
   unsigned x, y, z;
   __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {}
 };
 
 typedef struct cudaStream *cudaStream_t;
 
 int cudaConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0,
                       cudaStream_t stream = 0);
+#endif // !__NVCC__

test/SemaCUDA/attributes-on-non-cuda.cu

-// 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
 // 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}}

test/SemaCUDA/bad-attributes.cu

+// 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}}