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