Missing TargetPrefix for NVVM intrinsics

[Justin Lebar via llvm-commits]

> diff --git a/lib/Target/NVPTX/NVPTXISelLowering.cpp b/lib/Target/NVPTX/NVPTXISelLowering.cpp

Why are we losing the comment here?

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

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