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

[Richard Smith]

================
Comment at: lib/CodeGen/TargetInfo.cpp:5082-5083
@@ -5081,11 +5081,4 @@
       // 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) {
-        // Create !{<func-ref>, metadata !"minctasm", i32 <val>} node
-        addNVVMMetadata(F, "minctasm", MinCTASM);
+      llvm::APSInt MaxThreads;
+      if (Attr->getMaxThreads()->EvaluateAsInt(MaxThreads, M.getContext())) {
+        if (MaxThreads > 0)
----------------
You can use `Expr::EvaluateKnownConstInt` here.

================
Comment at: lib/CodeGen/TargetInfo.cpp:5085
@@ +5084,3 @@
+        if (MaxThreads > 0)
+          addNVVMMetadata(F, "maxntidx", MaxThreads.getExtValue());
+      } else
----------------
`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.

================
Comment at: lib/CodeGen/TargetInfo.cpp:5094
@@ +5093,3 @@
+        llvm::APSInt MinBlocks;
+        if (Attr->getMinBlocks()->EvaluateAsInt(MinBlocks, M.getContext())) {
+          if (MinBlocks > 0)
----------------
Likewise.

================
Comment at: lib/Sema/SemaDeclAttr.cpp:3490
@@ +3489,3 @@
+  // template arguments.
+  if (!E->getType()->isIntegerType() || E->containsUnexpandedParameterPack()) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_n_type)
----------------
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.)

================
Comment at: lib/Sema/SemaDeclAttr.cpp:3490
@@ +3489,3 @@
+  // template arguments.
+  if (!E->getType()->isIntegerType() || E->containsUnexpandedParameterPack()) {
+    S.Diag(Attr.getLoc(), diag::err_attribute_argument_n_type)
----------------
rsmith wrote:
> 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.)
If `E->getType()` is a dependent type, we presumably shouldn't reject here.

================
Comment at: lib/Sema/SemaTemplateInstantiateDecl.cpp:210
@@ +209,3 @@
+static Expr *
+getLaunchBoundsArgument(Sema &S,
+                        const MultiLevelTemplateArgumentList &TemplateArgs,
----------------
Can you refactor this to reuse the code from SemaDeclAttr rather than duplicating parts of it?

================
Comment at: lib/Sema/SemaTemplateInstantiateDecl.cpp:221-223
@@ +220,5 @@
+  if (ExprResult.isInvalid()) {
+    S.Diag(Attr.getLocation(), diag::err_attribute_argument_n_type)
+        << &Attr << Idx << AANT_ArgumentIntegerConstant
+        << TemplateArgExpr->getSourceRange();
+    return nullptr;
----------------
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.

http://reviews.llvm.org/D8985

EMAIL PREFERENCES
  http://reviews.llvm.org/settings/panel/emailpreferences/