[nvptx] Use nvptx_warpsync / nvptx_uniform_warp_check for -muniform-simt
authorTom de Vries <tdevries@suse.de>
Fri, 18 Feb 2022 15:50:03 +0000 (16:50 +0100)
committerTom de Vries <tdevries@suse.de>
Sat, 19 Feb 2022 18:57:12 +0000 (19:57 +0100)
commit8e5c34ab45f34aadea65c5ba33ec685264b6ec66
tree7b409182fd796d391d64e4617cec639805ab5934
parent9e3bbb4a8024121eb0fa675cb1f074218c1345a6
[nvptx] Use nvptx_warpsync / nvptx_uniform_warp_check for -muniform-simt

With the default ptx isa 6.0, we have for uniform-simt-1.c:
...
        @%r33   atom.global.cas.b32     %r26, [a], %r28, %r29;
                shfl.sync.idx.b32       %r26, %r26, %r32, 31, 0xffffffff;
...

The atomic insn is predicated by -muniform-simt, and the subsequent insn does
a warp sync, at which point the warp is uniform again.

But with -mptx=3.1, we have instead:
...
        @%r33   atom.global.cas.b32     %r26, [a], %r28, %r29;
                shfl.idx.b32    %r26, %r26, %r32, 31;
...

The shfl does not sync the warp, and we want the warp to go back to executing
uniformly asap.  We cannot enforce this, but at least check this using
nvptx_uniform_warp_check, similar to how that is done for openacc.

Likewise, detect the case that no shfl insn is emitted, and add a
nvptx_uniform_warp_check or nvptx_warpsync.

gcc/ChangeLog:

2022-02-19  Tom de Vries  <tdevries@suse.de>

* config/nvptx/nvptx.cc (nvptx_unisimt_handle_set): Change return
type to bool.
(nvptx_reorg_uniform_simt): Insert nvptx_uniform_warp_check or
nvptx_warpsync, if necessary.

gcc/testsuite/ChangeLog:

2022-02-19  Tom de Vries  <tdevries@suse.de>

* gcc.target/nvptx/uniform-simt-1.c: Add scan-assembler test.
* gcc.target/nvptx/uniform-simt-2.c: New test.
gcc/config/nvptx/nvptx.cc
gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c
gcc/testsuite/gcc.target/nvptx/uniform-simt-2.c [new file with mode: 0644]