[CUDA] Allow using integral non-type template parameters as launch_bounds attribute arguments.

[tra]

Summary

Allow using integral non-type template parameters as launch_bounds attribute arguments.

• Changed CUDALaunchBounds arguments from integers to Expr* so they can be saved in AST for instantiation, if needed.
• Added support for template instantiation of launch_bounds attrubute.
• Moved evaluation of launch_bounds arguments to NVPTXTargetCodeGenInfo:: SetTargetAttributes() where it can be done after template instantiation.
• Amended test cases.

[eliben]

The change looks fine. I'm not an expert on the template instantiation code, so you'll need to get an OK from Richard or someone else more familiar with this part.

lib/CodeGen/TargetInfo.cpp

5082 ↗

(On Diff #23630)

Can you just fold the FD->getAttr into the if () ? [it returns nullptr if the attr isn't there, I think]

5089 ↗

(On Diff #23630)

This comment was there before, but ISTM that s/default/optional/ makes more sense. What do you think?

lib/Sema/SemaTemplateInstantiateDecl.cpp

205 ↗

(On Diff #23630)

Document this function - what the arguments are, etc.

269 ↗

(On Diff #23630)

Fold the dyn_cast into the if condition here?

test/SemaCUDA/launch_bounds.cu

13 ↗

(On Diff #23630)

Can you add tests that use actual expressions?

Also a test that catches an invalid expression and complains as expected?

[tra]

Addressed review comments by eliben@ and aaron@

• Improved argument type/value checking.
• Added a warning on negative launch_bounds arguments and disabled corresponding PTX directive emission when that happens.
• Added test cases for various invalid launch_bounds arguments.

[tra]

lib/CodeGen/TargetInfo.cpp

5082 ↗

(On Diff #23630)

Done.

5089 ↗

(On Diff #23630)

Indeed. Fixed.

lib/Sema/SemaTemplateInstantiateDecl.cpp

205 ↗

(On Diff #23630)

Done.

269 ↗

(On Diff #23630)

OK.

test/SemaCUDA/launch_bounds.cu

13 ↗

(On Diff #23630)

Done and done.

[rsmith]

lib/CodeGen/TargetInfo.cpp

5082–5083 ↗

(On Diff #23709)

You can use Expr::EvaluateKnownConstInt here.

5085 ↗

(On Diff #23709)

getExtValue will assert if the int is wider than 64 bits. You should either add arbitrary width iN metadata here or put a range check on the value in Sema.

5094 ↗

(On Diff #23709)

Likewise.

lib/Sema/SemaDeclAttr.cpp

3490 ↗

(On Diff #23709)

It's better to use Sema::DiagnoseUnexpandedParameterPack to diagnose unexpanded packs. (In particular, it's sometimes valid to have an unexpanded pack here, if we're nested within a lambda-expression that is itself within a pack expansion, and DiagnoseUnexpandedParameterPack handles that case.)

3490 ↗

(On Diff #23709)

If E->getType() is a dependent type, we presumably shouldn't reject here.

lib/Sema/SemaTemplateInstantiateDecl.cpp

210 ↗

(On Diff #23709)

Can you refactor this to reuse the code from SemaDeclAttr rather than duplicating parts of it?

221–223 ↗

(On Diff #23709)

As a general rule, if someone hands you an invalid ExprResult, it means they've already produced a diagnostic and you should not do so.

[tra]

Addressed Richard Smith's review comments:

• Consolidated attribute creation and argument checking into Sema::AddLaunchBoundsAttr() which is now used from both SemaDeclAttr.cpp and SemaTemplateInstantiateDecl.cpp
• Let Expr::isIntegerConstantExpr() do the job of checking and reporting argument type errors.
• Use DiagnoseUnexpandedParameterPack() to report parameter packs when we reject them.

[tra]

lib/CodeGen/TargetInfo.cpp

5082–5083 ↗

(On Diff #23709)

Done.

5085 ↗

(On Diff #23709)

Added code in Sema to explicitly check whether result fits in 32 bits.

5094 ↗

(On Diff #23709)

Done.

lib/Sema/SemaDeclAttr.cpp

3490 ↗

(On Diff #23709)

Removed. isIntegerConstantExpr() does a pretty good job reporting errors in expressions it attempts to evaluate.

3490 ↗

(On Diff #23709)

Done.

lib/Sema/SemaTemplateInstantiateDecl.cpp

210 ↗

(On Diff #23709)

Done.

221–223 ↗

(On Diff #23709)

Fixed.

[rsmith]

LGTM with minor changes.

lib/Sema/SemaDeclAttr.cpp

3474–3475 ↗

(On Diff #23790)

There's no fundamental reason why value dependence must always imply instantiation dependence (and there are cases that should probably be modelled as being value-dependent but not instantiation-dependent, per discussion in the C++ committee, but we don't do so yet). Please also check for value-dependence when you check for instantiation-dependence.

3507 ↗

(On Diff #23790)

This comment seems out-of-place since it applies to dependent and non-dependent cases.

lib/Sema/SemaTemplateInstantiateDecl.cpp

213 ↗

(On Diff #23790)

This declaration, while correct, looks suspicious because only the second variable is initialized. Maybe declare MaxThreads below at the point where you assign to it, and declare MinBlocks immediately before the if (Attr.getMinBlocks()) line?

267 ↗

(On Diff #23790)

Add a newline between these, please.

[tra]

Few more changes to address Richard's comments.

[tra]

lib/Sema/SemaDeclAttr.cpp

3474–3475 ↗

(On Diff #23790)

Considering that we only care whether we can evaluate the expression or not, I've replaced isInstantiationDependent() with isValueDependent() and removed the assert.

3507 ↗

(On Diff #23790)

Removed.

lib/Sema/SemaTemplateInstantiateDecl.cpp

213 ↗

(On Diff #23790)

OK.

267 ↗

(On Diff #23790)

Done.

[rsmith]

LGTM

cfe/trunk/include/clang/Basic/Attr.td

def CUDAInvalidTarget : InheritableAttr {
   let Spellings = [];
   let Subjects = SubjectList<[Function]>;
   let LangOpts = [CUDA];
   let Documentation = [Undocumented];
 }
 
 def CUDALaunchBounds : InheritableAttr {
   let Spellings = [GNU<"launch_bounds">];
-  let Args = [IntArgument<"MaxThreads">, DefaultIntArgument<"MinBlocks", 0>];
+  let Args = [ExprArgument<"MaxThreads">, ExprArgument<"MinBlocks", 1>];
   let LangOpts = [CUDA];
   let Subjects = SubjectList<[ObjCMethod, FunctionLike], WarnDiag,
                              "ExpectedFunctionOrMethod">;
   // 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.
   let Documentation = [Undocumented];
 }

cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td

 def warn_concatenated_nsarray_literal : Warning<
   "concatenated NSString literal for an NSArray expression - "
   "possibly missing a comma">,
   InGroup<ObjCStringConcatenation>;
 def note_objc_literal_comparison_isequal : Note<
   "use 'isEqual:' instead">;
 def err_attribute_argument_is_zero : Error<
   "%0 attribute must be greater than 0">;
+def warn_attribute_argument_n_negative : Warning<
+  "%0 attribute parameter %1 is negative and will be ignored">,
+  InGroup<CudaCompat>;
 def err_property_function_in_objc_container : Error<
   "use of Objective-C property in function nested in Objective-C "
   "container not supported, move function outside its container">;
 
 let CategoryName = "Cocoa API Issue" in {
 def warn_objc_redundant_literal_use : Warning<
   "using %0 with a literal is redundant">, InGroup<ObjCRedundantLiteralUse>;
 }

cfe/trunk/include/clang/Sema/Sema.h

   void AddAssumeAlignedAttr(SourceRange AttrRange, Decl *D, Expr *E, Expr *OE,
                             unsigned SpellingListIndex);
 
   /// AddAlignValueAttr - Adds an align_value attribute to a particular
   /// declaration.
   void AddAlignValueAttr(SourceRange AttrRange, Decl *D, Expr *E,
                          unsigned SpellingListIndex);
 
+  /// AddLaunchBoundsAttr - Adds a launch_bounds attribute to a particular
+  /// declaration.
+  void AddLaunchBoundsAttr(SourceRange AttrRange, Decl *D, Expr *MaxThreads,
+                           Expr *MinBlocks, unsigned SpellingListIndex);
+
   // OpenMP directives and clauses.
 private:
   void *VarDataSharingAttributesStack;
   /// \brief Initialization of data-sharing attributes stack.
   void InitDataSharingAttributesStack();
   void DestroyDataSharingAttributesStack();
   ExprResult VerifyPositiveIntegerConstantInClause(Expr *Op,
                                                    OpenMPClauseKind CKind);

cfe/trunk/lib/CodeGen/TargetInfo.cpp

SetTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
   if (M.getLangOpts().CUDA) {
     // CUDA __global__ functions get a kernel metadata entry.  Since
     // __global__ functions cannot be called from the device, we do not
     // need to set the noinline attribute.
     if (FD->hasAttr<CUDAGlobalAttr>()) {
       // Create !{<func-ref>, metadata !"kernel", i32 1} node
       addNVVMMetadata(F, "kernel", 1);
     }
-    if (FD->hasAttr<CUDALaunchBoundsAttr>()) {
+    if (CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>()) {
       // Create !{<func-ref>, metadata !"maxntidx", i32 <val>} node
-      addNVVMMetadata(F, "maxntidx",
-                      FD->getAttr<CUDALaunchBoundsAttr>()->getMaxThreads());
-      // min blocks is a default argument for CUDALaunchBoundsAttr, so getting a
-      // zero value from getMinBlocks either means it was not specified in
-      // __launch_bounds__ or the user specified a 0 value. In both cases, we
-      // don't have to add a PTX directive.
-      int MinCTASM = FD->getAttr<CUDALaunchBoundsAttr>()->getMinBlocks();
-      if (MinCTASM > 0) {
+      llvm::APSInt MaxThreads(32);
+      MaxThreads = Attr->getMaxThreads()->EvaluateKnownConstInt(M.getContext());
+      if (MaxThreads > 0)
+        addNVVMMetadata(F, "maxntidx", MaxThreads.getExtValue());
+
+      // min blocks is an optional argument for CUDALaunchBoundsAttr. If it was
+      // not specified in __launch_bounds__ or if the user specified a 0 value,
+      // we don't have to add a PTX directive.
+      if (Attr->getMinBlocks()) {
+        llvm::APSInt MinBlocks(32);
+        MinBlocks = Attr->getMinBlocks()->EvaluateKnownConstInt(M.getContext());
+        if (MinBlocks > 0)
           // Create !{<func-ref>, metadata !"minctasm", i32 <val>} node
-        addNVVMMetadata(F, "minctasm", MinCTASM);
+          addNVVMMetadata(F, "minctasm", MinBlocks.getExtValue());
       }
     }
   }
 }
 
 void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::Function *F, StringRef Name,
                                              int Operand) {
   llvm::Module *M = F->getParent();

cfe/trunk/lib/Sema/SemaDeclAttr.cpp

Diag(Attr.getLoc(), diag::err_attribute_regparm_invalid_number)
       << Context.getTargetInfo().getRegParmMax() << NumParamsExpr->getSourceRange();
     Attr.setInvalid();
     return true;
   }
 
   return false;
 }
 
+// Checks whether an argument of launch_bounds attribute is acceptable
+// May output an error.
+static bool checkLaunchBoundsArgument(Sema &S, Expr *E,
+                                      const CUDALaunchBoundsAttr &Attr,
+                                      const unsigned Idx) {
+
+  if (S.DiagnoseUnexpandedParameterPack(E))
+    return false;
+
+  // Accept template arguments for now as they depend on something else.
+  // We'll get to check them when they eventually get instantiated.
+  if (E->isValueDependent())
+    return true;
+
+  llvm::APSInt I(64);
+  if (!E->isIntegerConstantExpr(I, S.Context)) {
+    S.Diag(E->getExprLoc(), diag::err_attribute_argument_n_type)
+        << &Attr << Idx << AANT_ArgumentIntegerConstant << E->getSourceRange();
+    return false;
+  }
+  // Make sure we can fit it in 32 bits.
+  if (!I.isIntN(32)) {
+    S.Diag(E->getExprLoc(), diag::err_ice_too_large) << I.toString(10, false)
+                                                     << 32 << /* Unsigned */ 1;
+    return false;
+  }
+  if (I < 0)
+    S.Diag(E->getExprLoc(), diag::warn_attribute_argument_n_negative)
+        << &Attr << Idx << E->getSourceRange();
+
+  return true;
+}
+
+void Sema::AddLaunchBoundsAttr(SourceRange AttrRange, Decl *D, Expr *MaxThreads,
+                               Expr *MinBlocks, unsigned SpellingListIndex) {
+  CUDALaunchBoundsAttr TmpAttr(AttrRange, Context, MaxThreads, MinBlocks,
+                               SpellingListIndex);
+
+  if (!checkLaunchBoundsArgument(*this, MaxThreads, TmpAttr, 0))
+    return;
+
+  if (MinBlocks && !checkLaunchBoundsArgument(*this, MinBlocks, TmpAttr, 1))
+    return;
+
+  D->addAttr(::new (Context) CUDALaunchBoundsAttr(
+      AttrRange, Context, MaxThreads, MinBlocks, SpellingListIndex));
+}
+
 static void handleLaunchBoundsAttr(Sema &S, Decl *D,
                                    const AttributeList &Attr) {
-  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))
+  if (!checkAttributeAtLeastNumArgs(S, Attr, 1) ||
+      !checkAttributeAtMostNumArgs(S, Attr, 2))
     return;
 
-  D->addAttr(::new (S.Context)
-              CUDALaunchBoundsAttr(Attr.getRange(), S.Context,
-                                  MaxThreads, MinBlocks,
-                                  Attr.getAttributeSpellingListIndex()));
+  S.AddLaunchBoundsAttr(Attr.getRange(), D, Attr.getArgAsExpr(0),
+                        Attr.getNumArgs() > 1 ? Attr.getArgAsExpr(1) : nullptr,
+                        Attr.getAttributeSpellingListIndex());
 }
 
 static void handleArgumentWithTypeTagAttr(Sema &S, Decl *D,
                                           const AttributeList &Attr) {
   if (!Attr.isArgIdent(0)) {
     S.Diag(Attr.getLoc(), diag::err_attribute_argument_n_type)
       << Attr.getName() << /* arg num = */ 1 << AANT_ArgumentIdentifier;
     return;

cfe/trunk/lib/Sema/SemaTemplateInstantiateDecl.cpp

static void instantiateDependentEnableIfAttr(
 
   EnableIfAttr *EIA = new (S.getASTContext())
                         EnableIfAttr(A->getLocation(), S.getASTContext(), Cond,
                                      A->getMessage(),
                                      A->getSpellingListIndex());
   New->addAttr(EIA);
 }
 
+// Constructs and adds to New a new instance of CUDALaunchBoundsAttr using
+// template A as the base and arguments from TemplateArgs.
+static void instantiateDependentCUDALaunchBoundsAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
+    const CUDALaunchBoundsAttr &Attr, Decl *New) {
+  // The alignment expression is a constant expression.
+  EnterExpressionEvaluationContext Unevaluated(S, Sema::ConstantEvaluated);
+
+  ExprResult Result = S.SubstExpr(Attr.getMaxThreads(), TemplateArgs);
+  if (Result.isInvalid())
+    return;
+  Expr *MaxThreads = Result.getAs<Expr>();
+
+  Expr *MinBlocks = nullptr;
+  if (Attr.getMinBlocks()) {
+    Result = S.SubstExpr(Attr.getMinBlocks(), TemplateArgs);
+    if (Result.isInvalid())
+      return;
+    MinBlocks = Result.getAs<Expr>();
+  }
+
+  S.AddLaunchBoundsAttr(Attr.getLocation(), New, MaxThreads, MinBlocks,
+                        Attr.getSpellingListIndex());
+}
+
 void Sema::InstantiateAttrs(const MultiLevelTemplateArgumentList &TemplateArgs,
                             const Decl *Tmpl, Decl *New,
                             LateInstantiatedAttrVec *LateAttrs,
                             LocalInstantiationScope *OuterMostScope) {
   for (const auto *TmplAttr : Tmpl->attrs()) {
     // FIXME: This should be generalized to more than just the AlignedAttr.
     const AlignedAttr *Aligned = dyn_cast<AlignedAttr>(TmplAttr);
     if (Aligned && Aligned->isAlignmentDependent()) {
for (const auto *TmplAttr : Tmpl->attrs()) {
 
     const EnableIfAttr *EnableIf = dyn_cast<EnableIfAttr>(TmplAttr);
     if (EnableIf && EnableIf->getCond()->isValueDependent()) {
       instantiateDependentEnableIfAttr(*this, TemplateArgs, EnableIf, Tmpl,
                                        New);
       continue;
     }
 
+    if (const CUDALaunchBoundsAttr *CUDALaunchBounds =
+            dyn_cast<CUDALaunchBoundsAttr>(TmplAttr)) {
+      instantiateDependentCUDALaunchBoundsAttr(*this, TemplateArgs,
+                                               *CUDALaunchBounds, New);
+      continue;
+    }
+
     // Existing DLL attribute on the instantiation takes precedence.
     if (TmplAttr->getKind() == attr::DLLExport ||
         TmplAttr->getKind() == attr::DLLImport) {
       if (New->hasAttr<DLLExportAttr>() || New->hasAttr<DLLImportAttr>()) {
         continue;
       }
     }
 

cfe/trunk/test/CodeGenCUDA/launch-bounds.cu

 __global__ void
 __launch_bounds__( MAX_THREADS_PER_BLOCK )
 Kernel2()
 {
 }
 }
 
 // CHECK: !{{[0-9]+}} = !{void ()* @Kernel2, !"maxntidx", i32 256}
+
+template <int max_threads_per_block>
+__global__ void
+__launch_bounds__(max_threads_per_block)
+Kernel3()
+{
+}
+
+template void Kernel3<MAX_THREADS_PER_BLOCK>();
+// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel3{{.*}}, !"maxntidx", i32 256}
+
+template <int max_threads_per_block, int min_blocks_per_mp>
+__global__ void
+__launch_bounds__(max_threads_per_block, min_blocks_per_mp)
+Kernel4()
+{
+}
+template void Kernel4<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();
+
+// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel4{{.*}}, !"maxntidx", i32 256}
+// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel4{{.*}}, !"minctasm", i32 2}
+
+const int constint = 100;
+template <int max_threads_per_block, int min_blocks_per_mp>
+__global__ void
+__launch_bounds__(max_threads_per_block + constint,
+                  min_blocks_per_mp + max_threads_per_block)
+Kernel5()
+{
+}
+template void Kernel5<MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP>();
+
+// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel5{{.*}}, !"maxntidx", i32 356}
+// CHECK: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel5{{.*}}, !"minctasm", i32 258}
+
+// Make sure we don't emit negative launch bounds values.
+__global__ void
+__launch_bounds__( -MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP )
+Kernel6()
+{
+}
+// CHECK-NOT: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel6{{.*}}, !"maxntidx",
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel6{{.*}}, !"minctasm",
+
+__global__ void
+__launch_bounds__( MAX_THREADS_PER_BLOCK, -MIN_BLOCKS_PER_MP )
+Kernel7()
+{
+}
+// CHECK:     !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"maxntidx",
+// CHECK-NOT: !{{[0-9]+}} = !{void ()* @{{.*}}Kernel7{{.*}}, !"minctasm",

cfe/trunk/test/SemaCUDA/launch_bounds.cu

-// RUN: %clang_cc1 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s
 
 #include "Inputs/cuda.h"
 
-__launch_bounds__(128, 7) void Test1(void);
-__launch_bounds__(128) void Test2(void);
+__launch_bounds__(128, 7) void Test2Args(void);
+__launch_bounds__(128) void Test1Arg(void);
 
-__launch_bounds__(1, 2, 3) void Test3(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
-__launch_bounds__() void Test4(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
+__launch_bounds__(0xffffffff) void TestMaxArg(void);
+__launch_bounds__(0x100000000) void TestTooBigArg(void); // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
+__launch_bounds__(0x10000000000000000) void TestWayTooBigArg(void); // expected-error {{integer literal is too large to be represented in any integer type}}
 
-int Test5 __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
+__launch_bounds__(-128, 7) void TestNegArg1(void); // expected-warning {{'launch_bounds' attribute parameter 0 is negative and will be ignored}}
+__launch_bounds__(128, -7) void TestNegArg2(void); // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
+
+__launch_bounds__(1, 2, 3) void Test3Args(void); // expected-error {{'launch_bounds' attribute takes no more than 2 arguments}}
+__launch_bounds__() void TestNoArgs(void); // expected-error {{'launch_bounds' attribute takes at least 1 argument}}
+
+int TestNoFunction __launch_bounds__(128, 7); // expected-warning {{'launch_bounds' attribute only applies to functions and methods}}
+
+__launch_bounds__(true) void TestBool(void);
+__launch_bounds__(128.0) void TestFP(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+__launch_bounds__((void*)0) void TestNullptr(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+
+int nonconstint = 256;
+__launch_bounds__(nonconstint) void TestNonConstInt(void); // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
+
+const int constint = 512;
+__launch_bounds__(constint) void TestConstInt(void);
+__launch_bounds__(constint * 2 + 3) void TestConstIntExpr(void);
+
+template <int a, int b> __launch_bounds__(a, b) void TestTemplate2Args(void) {}
+template void TestTemplate2Args<128,7>(void);
+
+template <int a> __launch_bounds__(a) void TestTemplate1Arg(void) {}
+template void TestTemplate1Arg<128>(void);
+
+template <class a>
+__launch_bounds__(a) void TestTemplate1ArgClass(void) {} // expected-error {{'a' does not refer to a value}}
+// expected-note@-2 {{declared here}}
+
+template <int a, int b, int c>
+__launch_bounds__(a + b, c + constint) void TestTemplateExpr(void) {}
+template void TestTemplateExpr<128+constint, 3, 7>(void);
+
+template <int... Args>
+__launch_bounds__(Args) void TestTemplateVariadicArgs(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}
+
+template <int... Args>
+__launch_bounds__(1, Args) void TestTemplateVariadicArgs2(void) {} // expected-error {{expression contains unexpanded parameter pack 'Args'}}