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

[Manjunath Kudlur]

================
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);
----------------
Eli Bendersky wrote:
> 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.
Done.

================
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);
----------------
Eli Bendersky wrote:
> Put this comment above the add* call above too (OpenCL case)?
Done.

================
Comment at: lib/CodeGen/TargetInfo.cpp:4844
@@ +4843,3 @@
+      // Create !{<func-ref>, metadata !"maxntidx", i32 <val>} node
+      addNVPTXMetadata(F, "maxntidx",
+                       FD->getAttr<CUDALaunchBoundsAttr>()->getMaxThreads());
----------------
Eli Bendersky wrote:
> Is it semantically correct to just populate the "x" here, rather than for all dimensions?
The code in lib/Target/NVPTX/NVPTXAsmPrinter.cpp populates the other values to default (1). So, yes.

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

================
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
----------------
Eli Bendersky wrote:
> Explain the default=0 case in a comment here
Done.


http://reviews.llvm.org/D3318