This is an archive of the discontinued LLVM Phabricator instance.

Added support for CUDA __launch_bounds__ attribute to CodeGen.
ClosedPublic

Authored by keveman on Apr 8 2014, 10:18 AM.

Details

Summary

Added support for CUDA launch_bounds attribute to CodeGen.
Sema does have a CUDALaunchBoundsAttr, but CodeGen was doing nothing with it.
This change translates CUDALaunchBoundsAttr to maxntidx and minctasm
metadata, which NVPTX then translates to the correct PTX directives.

Diff Detail

Event Timeline

eliben added inline comments.Apr 8 2014, 12:39 PM
lib/CodeGen/TargetInfo.cpp
4767

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.
4839

Put this comment above the add* call above too (OpenCL case)?

4844

Is it semantically correct to just populate the "x" here, rather than for all dimensions?

4846

no need for const?

4847

Explain the default=0 case in a comment here

keveman added inline comments.Apr 8 2014, 4:27 PM
lib/CodeGen/TargetInfo.cpp
4767

Done.

4839

Done.

4844

The code in lib/Target/NVPTX/NVPTXAsmPrinter.cpp populates the other values to default (1). So, yes.

4846

Done.

4847

Done.

keveman updated this revision to Unknown Object (????).Apr 8 2014, 4:33 PM

Updated according to review comments.

eliben accepted this revision.Apr 9 2014, 5:55 AM

LGTM

Justin, can you please take a look?

Trying one more time to see if the request for review reaches cfe-commits@.

rnk accepted this revision.Apr 14 2014, 5:15 PM

lgtm

lib/CodeGen/TargetInfo.cpp
4770

Having const ints feels silly, since the caller really doesn't care.

4845

nit: indentation

4850

The naming for local variables is StudlyCaps.

keveman updated this revision to Unknown Object (????).Apr 14 2014, 5:41 PM

Addressed rnk@'s comments.

lib/CodeGen/TargetInfo.cpp
4770

const removed.

4845

Fixed.

4850

Fixed.

Committed in r206302, and the test in r206303.

eliben closed this revision.Jun 20 2014, 7:13 AM