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

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.

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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20200313/870a4280/attachment.html>

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>