[PATCH] Added support for CUDA __launch_bounds__ attribute to CodeGen.

Eli Bendersky eliben at google.com
Tue Apr 8 12:39:45 PDT 2014



================
Comment at: lib/CodeGen/TargetInfo.cpp:4767
@@ -4766,2 +4766,3 @@
 private:
-  static void addKernelMetadata(llvm::Function *F);
+  static void addNVPTXMetadata(llvm::Function *F, StringRef Name,
+                               const int Operand);
----------------
Couple of notes:

1. Perhaps this is better named "addNVVMMetadata"? After all, the annotations are called nvvm.* not nvptx.*? 
2. Add documentation comment to this method explaining its arguments.

================
Comment at: lib/CodeGen/TargetInfo.cpp:4839
@@ +4838,3 @@
+    if (FD->hasAttr<CUDAGlobalAttr>()) {
+      // Create !{<func-ref>, metadata !"kernel", i32 1} node
+      addNVPTXMetadata(F, "kernel", 1);
----------------
Put this comment above the add* call above too (OpenCL case)?

================
Comment at: lib/CodeGen/TargetInfo.cpp:4844
@@ +4843,3 @@
+      // Create !{<func-ref>, metadata !"maxntidx", i32 <val>} node
+      addNVPTXMetadata(F, "maxntidx",
+                       FD->getAttr<CUDALaunchBoundsAttr>()->getMaxThreads());
----------------
Is it semantically correct to just populate the "x" here, rather than for all dimensions?

================
Comment at: lib/CodeGen/TargetInfo.cpp:4846
@@ +4845,3 @@
+                       FD->getAttr<CUDALaunchBoundsAttr>()->getMaxThreads());
+      const int minctasm = FD->getAttr<CUDALaunchBoundsAttr>()->getMinBlocks();
+      if (minctasm > 0) {
----------------
no need for const?

================
Comment at: lib/CodeGen/TargetInfo.cpp:4847
@@ +4846,3 @@
+      const int minctasm = FD->getAttr<CUDALaunchBoundsAttr>()->getMinBlocks();
+      if (minctasm > 0) {
+        // Create !{<func-ref>, metadata !"minctasm", i32 <val>} node
----------------
Explain the default=0 case in a comment here


http://reviews.llvm.org/D3318






More information about the cfe-commits mailing list