[PATCH] D35267: Pass Divergence Analysis data to selection DAG to drive divergence dependent instruction selection

[Stanislav Mekhanoshin via Phabricator via llvm-commits]

rampitec added a comment.

In https://reviews.llvm.org/D35267#950795, @alex-t wrote:

> In https://reviews.llvm.org/D35267#949981, @efriedma wrote:
>
> > > Any DAG transformation that change divergent pattern to not-divergent or vice versa is illegal.
> >
> > Transforming "x*0 -> 0" is illegal if x is divergent?  That seems surprising.
>
>
> Okay, I was unclear. Except for the constants. Your example is a corner case that turn the variable to the constant.
>  In this case w/o bit propagation we're still correct but sub-optimal.
>  I can imagine though the case where a long sequence of constant folding ends up with pure zero. If in addition the operand that becomes constant was the only divergent operand, we'd like to propagate.


More general there is a known corner case that for example "get_local_id(x) & ~63" is uniform.

The idea here is that get_local_id() is a source of divergence, but only low 6 bits of it are divergent and upper bits are uniform for our target. Handling such cases would need DA interface to be extended to produce a divergent bit mask instead of one bit answer and employ computeKnownBits on an expression to deduce expressions which converge to be uniform even if depend on a non-uniform value.

The corner case is order of magnitude less frequent than more straight forward uses of a divergent expression though. We have plans to extend divergence analysis in the future to handle this, but without a good mechanism to propagate through DAG it will not be very useful anyway.


https://reviews.llvm.org/D35267