[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

11–19

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

#include "Inputs/cuda.h"

33–34

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

33–34

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/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.
+//
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s
+
+// TODO: Check __cudart_builtin__, __nv_weak__,
+// __device_builtin_surface_type__, __device_builtin_texture_type__.
+
+// Make this file work with nvcc, for testing compatibility.
+#ifndef __NVCC__
+#define __device__ __attribute__((device))
+#define __host__ __attribute__((host))
+#define __constant__ __attribute__((constant))
+#define __shared__ __attribute__((shared))
+#define __global__ __attribute__((global))
+#define __device_builtin__ __attribute__((device_builtin))
+#endif
+
+// 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();
+__device_builtin__ int f1;
+__device_builtin__ void f2();
+
+// 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}}
+__device__ __device_builtin__ void z5();
+
+__host__ __device__ void z6();
+__host__ __global__ void z7();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__host__ __device_builtin__ void z8();
+
+__constant__ __device__ int z9;
+__constant__ __shared__ int z10;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__constant__ __device_builtin__ int z11;
+
+__shared__ __device__ int z12;
+__shared__ __constant__ int z13;  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__shared__ __device_builtin__ int z14;
+
+__global__ __device__ void z15();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__global__ __host__ void z16();  // expected-error {{attributes are not compatible}}
+// expected-note@-1 {{conflicting attribute is here}}
+__global__ __device_builtin__ void z17();