<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=iso-8859-1">
<style type="text/css" style="display:none;"> P {margin-top:0;margin-bottom:0;} </style>
</head>
<body dir="ltr">
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
I'm not sure what the usual "ping time" is for llvm-dev, but may I ask if there are any updates on this?</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
It appears that the following lines are the root cause of the reordering (<a href="https://github.com/llvm/llvm-project/blob/fdcb27105537f77c78c4473d4f7c47146ddbab69/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp#L2175" id="LPlnk679631">https://github.com/llvm/llvm-project/blob/fdcb27105537f77c78c4473d4f7c47146ddbab69/llvm/lib/Transforms/InstCombine/InstructionCombining.cpp#L2175</a>):</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<table>
<tbody>
<tr>
<td><br>
</td>
</tr>
<tr>
</tr>
</tbody>
</table>
<table>
<tbody>
<tr>
<td><span><span style="font-family: Consolas, Courier, monospace;">//</span></span><span style="font-family: Consolas, Courier, monospace;"> Handle gep(bitcast x) and gep(gep x, 0, 0, 0).</span></td>
</tr>
<span style="font-family: Consolas, Courier, monospace;"></span>
<tr>
</tr>
</tbody>
</table>
<table>
<tbody>
<tr>
<td><span style="font-family: Consolas, Courier, monospace;">Value *StrippedPtr = PtrOp-></span><span style="font-family: Consolas, Courier, monospace;">stripPointerCasts</span><span style="font-family: Consolas, Courier, monospace;">();</span></td>
</tr>
<span style="font-family: Consolas, Courier, monospace;"></span>
<tr>
</tr>
</tbody>
</table>
<table>
<tbody>
<tr>
<td><span style="font-family: Consolas, Courier, monospace;">PointerType *StrippedPtrTy = cast<PointerType>(StrippedPtr-></span><span style="font-family: Consolas, Courier, monospace;">getType</span><span style="font-family: Consolas, Courier, monospace;">());</span></td>
</tr>
<span style="font-family: Consolas, Courier, monospace;"></span>
<tr>
</tr>
</tbody>
</table>
<br>
</div>
<div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;">Namely, the </span>
<span style="font-family: Consolas, Courier, monospace;">stripPointerCasts</span><span style="font-family: Calibri, Helvetica, sans-serif;"> function also strips address spaces.</span></div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;">Would replacing this by
</span><span style="font-family: Consolas, Courier, monospace;">stripPointerCastsSameRepresentation</span><span style="font-family: Calibri, Helvetica, sans-serif;"> be an acceptable solution, or do other parts of LLVM rely on this reordering?</span></div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;">If not, I can send a patch in Phabricator to address this issue.</span></div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;"><br>
</span></div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;">Kind regards,</span></div>
<div><span style="font-family: Calibri, Helvetica, sans-serif;">Thomas Faingnaert<br>
</span></div>
<div style="font-family:Calibri,Arial,Helvetica,sans-serif; font-size:12pt; color:rgb(0,0,0);">
<br>
<hr tabindex="-1" style="display:inline-block; width:98%;">
<b>Van:</b> Johannes Doerfert<br>
<b>Verzonden:</b> Zaterdag, 14 Maart, 2020 6:50<br>
<b>Aan:</b> Thomas Faingnaert<br>
<b>CC:</b> llvm-dev@lists.llvm.org; Matthew Arsenault<br>
<b>Onderwerp:</b> Re: [llvm-dev] [InstCombine] Addrspacecast and GEP assumed commutative
<div><br>
</div>
</div>
<div class="BodyFragment"><font size="2"><span style="font-size:11pt;">
<div class="PlainText">On 03/13, Thomas Faingnaert via llvm-dev wrote:<br>
> I have a question regarding the InstCombine pass.<br>
> Given the following IR:<br>
> <br>
> define void @test(i64) {<br>
> %ptr = inttoptr i64 %0 to i16*<br>
> <br>
> %asc = addrspacecast i16* %ptr to i16 addrspace(1)*<br>
> %gep = getelementptr i16, i16 addrspace(1)* %asc, i64 16<br>
> store i16 0, i16 addrspace(1)* %gep, align 16<br>
> <br>
> ret void<br>
> }<br>
> <br>
> opt -instcombine transforms this to:<br>
> <br>
> define void @test(i64 %0) {<br>
> %ptr = inttoptr i64 %0 to i16*<br>
> <br>
> %gep1 = getelementptr i16, i16* %ptr, i64 16<br>
> %gep = addrspacecast i16* %gep1 to i16 addrspace(1)*<br>
> store i16 0, i16 addrspace(1)* %gep, align 16<br>
> <br>
> ret void<br>
> }<br>
> <br>
> As you can see, the order of the GEP and addrspacecast is changed.<br>
> 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.<br>
<br>
Yeah, I can see this being a problem. I think it is valid if the gep is<br>
inbounds but not otherwise. Matt, wdyt?<br>
<br>
<br>
> I stumbled upon this because this reordering actually gives problems for instruction selection in the NVPTX backend.<br>
> Without reordering, the above IR gets lowered to:<br>
> <br>
> ld.param.u64 %rd1, [test_param_0];<br>
> cvta.to.global.u64 %rd2, %rd1;<br>
> mov.u16 %rs1, 0;<br>
> st.global.u16 [%rd2+32], %rs1;<br>
> ret;<br>
> <br>
> But with the reordering, the backend instead emits this:<br>
> <br>
> ld.param.u64 %rd1, [test_param_0];<br>
> add.s64 %rd2, %rd1, 32;<br>
> cvta.to.global.u64 %rd3, %rd2;<br>
> mov.u16 %rs1, 0;<br>
> st.global.u16 [%rd3], %rs1;<br>
> ret;<br>
> <br>
> i.e. an explicit add instruction instead of folding the addition in the addressing mode of the store.<br>
> <br>
> My question is twofold:<br>
> <br>
> 1. Is this reordering of GEP and ASC in the InstCombine pass the expected behaviour?<br>
> 2. At the moment, I solved the issue described above by reordering the GEP and ASC during ISEL of NVPTX (<a href="https://reviews.llvm.org/D75817" target="_blank" rel="noopener noreferrer" data-auth="NotApplicable">https://reviews.llvm.org/D75817</a>),
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.<br>
<br>
> _______________________________________________<br>
> LLVM Developers mailing list<br>
> llvm-dev@lists.llvm.org<br>
> <a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" target="_blank" rel="noopener noreferrer" data-auth="NotApplicable">
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
<br>
<br>
-- <br>
<br>
Johannes Doerfert<br>
Researcher<br>
<br>
Argonne National Laboratory<br>
Lemont, IL 60439, USA<br>
<br>
jdoerfert@anl.gov<br>
</div>
</span></font></div>
</div>
</body>
</html>