From: Aaron Ballman <aaron@aaronballman.com>
Date: Thu, 19 Dec 2013 00:41:31 +0000 (+0000)
Subject: Added a comment about the launch_bounds attribute's AST node being required. Since... 
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=9be81e154324f38a9528bc3dce021c5fa07a5d69;p=clang

Added a comment about the launch_bounds attribute's AST node being required. Since there were no test cases for the attribute, some have been added. This promptly demonstrated a bug with the semantic handling, which is also fixed.

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 b8a60fb10a..8249965022 100644
--- 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 01a1bed0d4..376d75ce6f 100644
--- 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 26a8df0440..a9a4595a14 100644
--- 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
index 0000000000..c508a00c2d
--- /dev/null
+++ b/test/SemaCUDA/launch_bounds.cu
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -verify %s
+
+#include "cuda.h"
+
+__launch_bounds__(128, 7) void Test1(void);
+__launch_bounds__(128) void Test2(void);
+
+__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{attribute takes no more than 2 arguments}}
+
+// FIXME: the error should read that the attribute takes exactly one or two arguments, but there
+// is no support for such a diagnostic currently.
+__launch_bounds__() void Test4(void); // expected-error {{attribute takes no more than 2 arguments}}
+
+int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}