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.