[Bug target/97338] New: [nvptx] Convergence checking

[vries at gcc dot gnu.org]

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

            Bug ID: 97338
           Summary: [nvptx] Convergence checking
           Product: gcc
           Version: 11.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: ---

With ptx, we have insns that need to be executed in convergent mode, that is,
all threads in the warp active.

We can unfortunately not enforce this, but we could check it, which could help
pinpoint problems.

A ptx insn:
...
  vote.ballot.b32 %rbla, 1;
...
gives us the regmask of active threads, so we could check:
...
{
  .reg .u32 %rwarp_active_mask;
  vote.ballot.b32 %rwarp_active_mask, 1;
  .reg .pred %pconvergent;
  setp.eq.u32 %pconvergent,%rwarp_active_mask,-1;
  @ ! %pconvergent trap;
}
...