[PATCH] D116053: [MachineSink] Allow sinking of constant or ignorable physreg uses

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec added a comment.

In D116053#3205477 <https://reviews.llvm.org/D116053#3205477>, @vangthao wrote:

>> IR is essentially a single thread representation. The implicit exec use is our way to model mutithreaded divergence. Consider this transformation which shall now become legal:
>>
>>   int lid = get_local_id(0);      int lid = get_local_id(0);
>>   int i = 0;                      int i = 0;
>>   x = def();                      do {
>>   do {                        =>    x = def();
>>     use1(x);                        use1(x);
>>   } while(i++ < lid);             } while(i++ < lid);
>>   use2(x);                        use2(x);
>>
>> def dominates use2 in both cases, but in the second case not with every lane. All lanes except first will use an undef.
>
> We will not sink into a loop if the def is outside of the loop. In the test case `loop_sink_fmac`, the def was already in a loop and was why it was able to be sinked.

If you are sure then it is OK. Note my example was with a loop which is guaranteed to execute at least one iteration. @arsenm ?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D116053/new/

https://reviews.llvm.org/D116053