[Bug target/104991] New: [nvptx] Simplify muniform-simt transformation

[vries at gcc dot gnu.org]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=104991

            Bug ID: 104991
           Summary: [nvptx] Simplify muniform-simt transformation
           Product: gcc
           Version: 12.0
            Status: UNCONFIRMED
          Severity: enhancement
          Priority: P3
         Component: target
          Assignee: unassigned at gcc dot gnu.org
          Reporter: vries at gcc dot gnu.org
  Target Milestone: ---

The muniform-simt reorg pass transforms the insn stream, both inside and
outside an SIMT region.

The transform rewrites atomic insns by adding a predicate to execute in one
thread only outside the SIMT region.  Furthermore, if the atomic insn has a
result, a shuffle is added to propagate the result to all threads in the warp.

Inside the SIMT region, the predicate evaluates to true such that all threads
in the warp execute it.

And the source lane register for the shuffle is set such that the shuffle is a
nop inside the SIMT region.

However, since we've started using shfl.sync for the shuffle, the shuffle now
has it's own predicate, and consequently having a source lane register with
different values inside and outside the SIMT region is no longer necessary.