[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.