Added a comment about the launch_bounds attribute's AST node being required. Since...

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@197637 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td
index b8a60fb10a28f8f39d2799d7544ef781c9a11b31..82499650223030f6ab34ad6f98b1d758396f6b0b 100644 (file)
--- a/include/clang/Basic/Attr.td
+++ b/include/clang/Basic/Attr.td
@@ -414,6 +414,9 @@ def CUDALaunchBounds : InheritableAttr {
   let Spellings = [GNU<"launch_bounds">];
   let Args = [IntArgument<"MaxThreads">, DefaultIntArgument<"MinBlocks", 0>];
   let LangOpts = [CUDA];
+  // An AST node is created for this attribute, but is not used by other parts
+  // of the compiler. However, this node needs to exist in the AST because
+  // non-LLVM backends may be relying on the attribute's presence.
 }
 
 def CUDAShared : InheritableAttr {

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 01a1bed0d4158b4d14cf5bf0edaacc0c13457a15..376d75ce6f66b6a8c00ed818d736a174dbaa4ba2 100644 (file)
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -3352,7 +3352,8 @@ bool Sema::CheckRegparmAttr(const AttributeList &Attr, unsigned &numParams) {
   return false;
 }
 
-static void handleLaunchBoundsAttr(Sema &S, Decl *D, const AttributeList &Attr){
+static void handleLaunchBoundsAttr(Sema &S, Decl *D,
+                                   const AttributeList &Attr) {
   // check the attribute arguments.
   if (Attr.getNumArgs() != 1 && Attr.getNumArgs() != 2) {
     // FIXME: 0 is not okay.
@@ -3366,9 +3367,12 @@ static void handleLaunchBoundsAttr(Sema &S, Decl *D, const AttributeList &Attr){
     return;
   }
 
-  uint32_t MaxThreads, MinBlocks;
-  if (!checkUInt32Argument(S, Attr, Attr.getArgAsExpr(0), MaxThreads, 1) ||
-      !checkUInt32Argument(S, Attr, Attr.getArgAsExpr(1), MinBlocks, 2))
+  uint32_t MaxThreads, MinBlocks = 0;
+  if (!checkUInt32Argument(S, Attr, Attr.getArgAsExpr(0), MaxThreads, 1))
+    return;
+  if (Attr.getNumArgs() > 1 && !checkUInt32Argument(S, Attr,
+                                                    Attr.getArgAsExpr(1),
+                                                    MinBlocks, 2))
     return;
 
   D->addAttr(::new (S.Context)

diff --git a/test/SemaCUDA/cuda.h b/test/SemaCUDA/cuda.h
index 26a8df0440f151297f5d2ddce6be8673c86d76d7..a9a4595a14a96f18f1e73a2933f08454607da94b 100644 (file)
--- a/test/SemaCUDA/cuda.h
+++ b/test/SemaCUDA/cuda.h
@@ -7,6 +7,7 @@
 #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;

diff --git a/test/SemaCUDA/launch_bounds.cu b/test/SemaCUDA/launch_bounds.cu
new file mode 100644 (file)
index 0000000..c508a00
--- /dev/null
+++ b/test/SemaCUDA/launch_bounds.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s\r
+\r
+#include "cuda.h"\r
+\r
+__launch_bounds__(128, 7) void Test1(void);\r
+__launch_bounds__(128) void Test2(void);\r
+\r
+__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{attribute takes no more than 2 arguments}}\r
+\r
+// FIXME: the error should read that the attribute takes exactly one or two arguments, but there\r
+// is no support for such a diagnostic currently.\r
+__launch_bounds__() void Test4(void); // expected-error {{attribute takes no more than 2 arguments}}\r
+\r
+int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}\r