[llvm-dev] [InstCombine] Addrspacecast and GEP assumed commutative

[Johannes Doerfert via llvm-dev]

On 03/13, Thomas Faingnaert via llvm-dev wrote:
> I have a question regarding the InstCombine pass.
> Given the following IR:
> 
> define void @test(i64) {
>   %ptr = inttoptr i64 %0 to i16*
> 
>   %asc = addrspacecast i16* %ptr to i16 addrspace(1)*
>   %gep = getelementptr i16, i16 addrspace(1)* %asc, i64 16
>   store i16 0, i16 addrspace(1)* %gep, align 16
> 
>   ret void
> }
> 
> opt -instcombine transforms this to:
> 
> define void @test(i64 %0) {
>   %ptr = inttoptr i64 %0 to i16*
> 
>   %gep1 = getelementptr i16, i16* %ptr, i64 16
>   %gep = addrspacecast i16* %gep1 to i16 addrspace(1)*
>   store i16 0, i16 addrspace(1)* %gep, align 16
> 
>   ret void
> }
> 
> As you can see, the order of the GEP and addrspacecast is changed.
> Is this the expected behaviour? Normally, ASC(GEP(x)) is not necessarily equivalent to GEP(ASC(x)), e.g. when the GEP moves past the end of one address space.

Yeah, I can see this being a problem. I think it is valid if the gep is
inbounds but not otherwise. Matt, wdyt?

 
> I stumbled upon this because this reordering actually gives problems for instruction selection in the NVPTX backend.
> Without reordering, the above IR gets lowered to:
> 
> ld.param.u64 %rd1, [test_param_0];
> cvta.to.global.u64 %rd2, %rd1;
> mov.u16 %rs1, 0;
> st.global.u16 [%rd2+32], %rs1;
> ret;
> 
> But with the reordering, the backend instead emits this:
> 
> ld.param.u64 %rd1, [test_param_0];
> add.s64 %rd2, %rd1, 32;
> cvta.to.global.u64 %rd3, %rd2;
> mov.u16 %rs1, 0;
> st.global.u16 [%rd3], %rs1;
> ret;
> 
> i.e. an explicit add instruction instead of folding the addition in the addressing mode of the store.
> 
> My question is twofold:
> 
>   1.  Is this reordering of GEP and ASC in the InstCombine pass the expected behaviour?
>   2.  At the moment, I solved the issue described above by reordering the GEP and ASC during ISEL of NVPTX (https://reviews.llvm.org/D75817), but I don't check if the reordering is valid. If the reordering in InstCombine is indeed the expected behaviour, would it not be better to disable it? I imagine that verifying if the necessary conditions hold for reordering during ISEL will be quite complex.

> _______________________________________________
> LLVM Developers mailing list
> llvm-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev


-- 

Johannes Doerfert
Researcher

Argonne National Laboratory
Lemont, IL 60439, USA

jdoerfert at anl.gov
-------------- next part --------------
A non-text attachment was scrubbed...
Name: signature.asc
Type: application/pgp-signature
Size: 228 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20200314/006cce29/attachment.sig>