[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