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