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

[Thomas Faingnaert via llvm-dev]

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>