Missing TargetPrefix for NVVM intrinsics

Justin Bogner via llvm-commits llvm-commits at lists.llvm.org
Thu Jul 7 09:50:55 PDT 2016

Justin Lebar <jlebar at google.com> writes:
>> diff --git a/lib/Target/NVPTX/NVPTXISelLowering.cpp
>> b/lib/Target/NVPTX/NVPTXISelLowering.cpp
> Why are we losing the comment here?

That should be separate. AFAICT the comment is discussing intrinsics
that have never existed in open source and are not handled in
NVPTXTargetLowering::getTgtMemIntrinsic. I've left the comment there for

>> +// TODO: Why is this WARP_SZ rather than %warpsize, like the others?
> There's no %warpsize reg, afaict.  There is a %WARP_SZ, which is
> technically a register, while WARP_SZ, which is technically an integer
> constant.  I strongly suspect ptxas will compile these both to the
> same thing, but I don't see any harm in using the constant.
> http://docs.nvidia.com/cuda/parallel-thread-execution/#identifiers
> Otherwise looks good to me.

Committed in LLVM r274769 and clang r274770.

> On Wed, Jul 6, 2016 at 3:39 PM, Justin Bogner <mail at justinbogner.com> wrote:
>> Justin Lebar <jlebar at google.com> writes:
>>>> I'll wait on the ones that drop the ptx intrinsics and builtins
>>>> until I hear back about the ones that don't seem to have an nvvm
>>>> equivalent.
>>> What's the outstanding question, exactly?  It seems to me we should
>>> just rename them.
>> Ah, I misunderstood before.
>> This ended up being a bit more work on the LLVM side, but the attached
>> patch replaces all of the ptx.read intrinsics with nvvm.read.ptx.sreg
>> variants, cleans up the duplicate nvvm variants that leaves around, and
>> renames ptx.bar.sync to nvvm.bar.sync.
>> The clang patch just updates the names and doesn't change anything with
>> this version.
>> WDYT?

More information about the llvm-commits mailing list