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

[Eli Bendersky]

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.


================
Comment at: lib/CodeGen/TargetInfo.cpp:5082
@@ -5081,6 +5081,3 @@
       // 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
+      CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>();
+      llvm::APSInt MaxThreads;
----------------
Can you just fold the FD->getAttr into the if () ? [it returns nullptr if the attr isn't there, I think]

================
Comment at: lib/CodeGen/TargetInfo.cpp:5089
@@ +5088,3 @@
+
+      // min blocks is a default argument for CUDALaunchBoundsAttr. If it was
+      // not specified in __launch_bounds__ or the user specified a 0 value, we
----------------
This comment was there before, but ISTM that s/default/optional/ makes more sense. What do you think?

================
Comment at: lib/Sema/SemaTemplateInstantiateDecl.cpp:205
@@ -204,1 +204,3 @@
 
+static void instantiateDependentCUDALaunchBoundsAttr(
+    Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
----------------
Document this function - what the arguments are, etc.

================
Comment at: lib/Sema/SemaTemplateInstantiateDecl.cpp:269
@@ +268,3 @@
+        dyn_cast<CUDALaunchBoundsAttr>(TmplAttr);
+    if (CUDALaunchBounds) {
+      instantiateDependentCUDALaunchBoundsAttr(*this, TemplateArgs,
----------------
Fold the dyn_cast into the if condition here?

================
Comment at: test/SemaCUDA/launch_bounds.cu:13
@@ +12,3 @@
+
+template <int a, int b> __launch_bounds__(a, b) void Test6(void) {}
+template void Test6<128,7>(void);
----------------
Can you add tests that use actual expressions? 

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

http://reviews.llvm.org/D8985

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