Missing TargetPrefix for NVVM intrinsics

Justin Lebar via llvm-commits llvm-commits at lists.llvm.org
Wed Jul 6 16:31:35 PDT 2016

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


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.

More information about the llvm-commits mailing list