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

[Vang Thao via Phabricator via llvm-commits]

vangthao added a comment.

> 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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D116053