r197637 - 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.

[Aaron Ballman]

Author: aaronballman
Date: Wed Dec 18 18:41:31 2013
New Revision: 197637

URL: http://llvm.org/viewvc/llvm-project?rev=197637&view=rev
Log:
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.

Added:
    cfe/trunk/test/SemaCUDA/launch_bounds.cu
Modified:
    cfe/trunk/include/clang/Basic/Attr.td
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/SemaCUDA/cuda.h

Modified: cfe/trunk/include/clang/Basic/Attr.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=197637&r1=197636&r2=197637&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/Attr.td (original)
+++ cfe/trunk/include/clang/Basic/Attr.td Wed Dec 18 18:41:31 2013
@@ -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 {

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=197637&r1=197636&r2=197637&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Wed Dec 18 18:41:31 2013
@@ -3352,7 +3352,8 @@ bool Sema::CheckRegparmAttr(const Attrib
   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
     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)

Modified: cfe/trunk/test/SemaCUDA/cuda.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/cuda.h?rev=197637&r1=197636&r2=197637&view=diff
==============================================================================
--- cfe/trunk/test/SemaCUDA/cuda.h (original)
+++ cfe/trunk/test/SemaCUDA/cuda.h Wed Dec 18 18:41:31 2013
@@ -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;

Added: cfe/trunk/test/SemaCUDA/launch_bounds.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaCUDA/launch_bounds.cu?rev=197637&view=auto
==============================================================================
--- cfe/trunk/test/SemaCUDA/launch_bounds.cu (added)
+++ cfe/trunk/test/SemaCUDA/launch_bounds.cu Wed Dec 18 18:41:31 2013
@@ -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}}