<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 have a question regarding the InstCombine pass.</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Given the following IR:</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<pre><span style="font-family: Consolas, Courier, monospace;">define</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">void</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">@test</span><span style="font-family: Consolas, Courier, monospace;">(</span><span style="font-family: Consolas, Courier, monospace;">i64</span><span style="font-family: Consolas, Courier, monospace;">)</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">{</span><span style="font-family: Consolas, Courier, monospace;">
</span><span style="font-family: Consolas, Courier, monospace;">%ptr</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">=</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">inttoptr</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i64</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">%0</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">to</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;">*</span>
<span style="font-family: Consolas, Courier, monospace;">%asc</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">=</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">addrspacecast</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;">*</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">%ptr</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">to</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">addrspace</span><span style="font-family: Consolas, Courier, monospace;">(</span><span style="font-family: Consolas, Courier, monospace;">1</span><span style="font-family: Consolas, Courier, monospace;">)*</span><span style="font-family: Consolas, Courier, monospace;">
</span><span style="font-family: Consolas, Courier, monospace;">%gep</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">=</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">getelementptr</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;">,</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">addrspace</span><span style="font-family: Consolas, Courier, monospace;">(</span><span style="font-family: Consolas, Courier, monospace;">1</span><span style="font-family: Consolas, Courier, monospace;">)*</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">%asc</span><span style="font-family: Consolas, Courier, monospace;">,</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i64</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">16</span><span style="font-family: Consolas, Courier, monospace;">
</span><span style="font-family: Consolas, Courier, monospace;">store</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">0</span><span style="font-family: Consolas, Courier, monospace;">,</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">i16</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">addrspace</span><span style="font-family: Consolas, Courier, monospace;">(</span><span style="font-family: Consolas, Courier, monospace;">1</span><span style="font-family: Consolas, Courier, monospace;">)*</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">%gep</span><span style="font-family: Consolas, Courier, monospace;">,</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">align</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">16</span>
<span style="font-family: Consolas, Courier, monospace;">ret</span><span style="font-family: Consolas, Courier, monospace;"> </span><span style="font-family: Consolas, Courier, monospace;">void</span>
<span style="font-family: Consolas, Courier, monospace;">}</span></pre>
</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<span style="font-family: Consolas, Courier, monospace;">opt -instcombine</span> transforms this to:</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);">
<span style="font-family: Consolas, Courier, monospace;">define void @test(i64 %0) {</span><br>
<div><span style="font-family: Consolas, Courier, monospace;"> %ptr = inttoptr i64 %0 to i16*</span></div>
<div><span style="font-family: Consolas, Courier, monospace;"></span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;"> %gep1 = getelementptr i16, i16* %ptr, i64 16</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;"> %gep = addrspacecast i16* %gep1 to i16 addrspace(1)*</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;"> store i16 0, i16 addrspace(1)* %gep, align 16</span></div>
<div><span style="font-family: Consolas, Courier, monospace;"></span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;"> ret void</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">}</span><br>
</div>
</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);">
As you can see, the order of the GEP and addrspacecast is changed.</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
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.</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);">
I stumbled upon this because this reordering actually gives problems for instruction selection in the NVPTX backend.</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
Without reordering, the above IR gets lowered to:</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);">
<span style="font-family: Consolas, Courier, monospace;">ld.param.u64 %rd1, [test_param_0];</span><br>
<div><span style="font-family: Consolas, Courier, monospace;">cvta.to.global.u64 %rd2, %rd1;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">mov.u16 %rs1, 0;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">st.global.u16 [%rd2+32], %rs1;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">ret;</span><span></span></div>
</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);">
But with the reordering, the backend instead emits this:</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<br>
<div><span style="font-family: Consolas, Courier, monospace;">ld.param.u64 %rd1, [test_param_0];</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">add.s64 %rd2, %rd1, 32;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">cvta.to.global.u64 %rd3, %rd2;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">mov.u16 %rs1, 0;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">st.global.u16 [%rd3], %rs1;</span><br>
</div>
<div><span style="font-family: Consolas, Courier, monospace;">ret;</span><span></span></div>
</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);">
i.e. an explicit add instruction instead of folding the addition in the addressing mode of the store.</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);">
My question is twofold:</div>
<div style="font-family: Calibri, Helvetica, sans-serif; font-size: 12pt; color: rgb(0, 0, 0);">
<ol>
<li>Is this reordering of GEP and ASC in the InstCombine pass the expected behaviour?</li><li>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" id="LPlnk533040">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>
</li></ol>
</div>
</body>
</html>