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

Artem Belevich tra at google.com
Mon Apr 13 18:26:22 PDT 2015


================
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;
----------------
eliben wrote:
> Can you just fold the FD->getAttr into the if () ? [it returns nullptr if the attr isn't there, I think]
Done.

================
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
----------------
eliben wrote:
> This comment was there before, but ISTM that s/default/optional/ makes more sense. What do you think?
Indeed. Fixed.

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

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

================
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);
----------------
eliben wrote:
> Can you add tests that use actual expressions? 
> 
> Also a test that catches an invalid expression and complains as expected?
Done and done.

http://reviews.llvm.org/D8985

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






More information about the cfe-commits mailing list