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

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec added a comment.

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

> In D116053#3203538 <https://reviews.llvm.org/D116053#3203538>, @rampitec wrote:
>
>> I am not sure I can prove to myself this is legal. For example you are sinking a def into a loop with divergent condition and this def is used after the loop. Can this happen? If so a def might be done with an exec smaller than a use which creates an undef. Hoisting was OK because def was moved into a direction where exec is strictly not less than before. Did you run PSDB on it?
>
> I believe the requirement for defs to dominate all uses prevents this from happening but I can add more test cases to check for this. This passed PSDB.

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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D116053