Missing TargetPrefix for NVVM intrinsics

[Justin Bogner via llvm-commits]

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

>> +// 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?
>>