public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
* [Bug target/97338] New: [nvptx] Convergence checking
@ 2020-10-08 16:14 vries at gcc dot gnu.org
  2022-02-23 11:45 ` [Bug target/97338] " vries at gcc dot gnu.org
  0 siblings, 1 reply; 2+ messages in thread
From: vries at gcc dot gnu.org @ 2020-10-08 16:14 UTC (permalink / raw)
  To: gcc-bugs

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

^ permalink raw reply	[flat|nested] 2+ messages in thread

* [Bug target/97338] [nvptx] Convergence checking
  2020-10-08 16:14 [Bug target/97338] New: [nvptx] Convergence checking vries at gcc dot gnu.org
@ 2022-02-23 11:45 ` vries at gcc dot gnu.org
  0 siblings, 0 replies; 2+ messages in thread
From: vries at gcc dot gnu.org @ 2022-02-23 11:45 UTC (permalink / raw)
  To: gcc-bugs

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

Tom de Vries <vries at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
         Resolution|---                         |FIXED
   Target Milestone|---                         |12.0
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #1 from Tom de Vries <vries at gcc dot gnu.org> ---
openacc:
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=f32f74c2e8cef5fe37af6d4e8d7e8f6b4c8ae9a8

openmp:
https://gcc.gnu.org/git/?p=gcc.git;a=commit;h=8e5c34ab45f34aadea65c5ba33ec685264b6ec66

^ permalink raw reply	[flat|nested] 2+ messages in thread

end of thread, other threads:[~2022-02-23 11:45 UTC | newest]

Thread overview: 2+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2020-10-08 16:14 [Bug target/97338] New: [nvptx] Convergence checking vries at gcc dot gnu.org
2022-02-23 11:45 ` [Bug target/97338] " vries at gcc dot gnu.org

This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).