OpenE2K
/
gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[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.
This commit is contained in:
Tom de Vries 2022-02-18 16:50:03 +01:00
parent 9e3bbb4a80
commit 8e5c34ab45
3 changed files with 52 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

34
gcc/config/nvptx/nvptx.cc
View File

@ -3248,12 +3248,18 @@ nvptx_call_insn_is_syscall_p (rtx_insn *insn)
/* If SET subexpression of INSN sets a register, emit a shuffle instruction to
propagate its value from lane MASTER to current lane. */
static void
static bool
nvptx_unisimt_handle_set (rtx set, rtx_insn *insn, rtx master)
{
rtx reg;
if (GET_CODE (set) == SET && REG_P (reg = SET_DEST (set)))
emit_insn_after (nvptx_gen_shuffle (reg, reg, master, SHUFFLE_IDX), insn);
{
emit_insn_after (nvptx_gen_shuffle (reg, reg, master, SHUFFLE_IDX),
insn);
return true;
}
return false;
}
/* Adjust code for uniform-simt code generation variant by making atomics and
@ -3275,8 +3281,30 @@ nvptx_reorg_uniform_simt ()
continue;
rtx pat = PATTERN (insn);
rtx master = nvptx_get_unisimt_master ();
bool shuffle_p = false;
for (int i = 0; i < XVECLEN (pat, 0); i++)
nvptx_unisimt_handle_set (XVECEXP (pat, 0, i), insn, master);
shuffle_p
|= nvptx_unisimt_handle_set (XVECEXP (pat, 0, i), insn, master);
if (shuffle_p && TARGET_PTX_6_0)
{
/* The shuffle is a sync, so uniformity is guaranteed. */
}
else
{
if (TARGET_PTX_6_0)
{
gcc_assert (!shuffle_p);
/* Emit after the insn, to guarantee uniformity. */
emit_insn_after (gen_nvptx_warpsync (), insn);
}
else
{
/* Emit after the insn (and before the shuffle, if there are any)
to check uniformity. */
emit_insn_after (gen_nvptx_uniform_warp_check (), insn);
}
}
rtx pred = nvptx_get_unisimt_predicate ();
pred = gen_rtx_NE (BImode, pred, const0_rtx);
pat = gen_rtx_COND_EXEC (VOIDmode, pred, pat);

1
gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c
View File

@ -16,3 +16,4 @@ f (void)
}
/* { dg-final { scan-assembler-times "@%r\[0-9\]*\tatom.global.cas" 1 } } */
/* { dg-final { scan-assembler-times "shfl.sync.idx.b32" 1 } } */

20
gcc/testsuite/gcc.target/nvptx/uniform-simt-2.c Normal file
View File

@ -0,0 +1,20 @@
/* { dg-options "-O2 -muniform-simt -mptx=3.1" } */
enum memmodel
{
MEMMODEL_RELAXED = 0,
};
int a = 0;
int
f (void)
{
int expected = 1;
return __atomic_compare_exchange_n (&a, &expected, 0, 0, MEMMODEL_RELAXED,
MEMMODEL_RELAXED);
}
/* { dg-final { scan-assembler-times "@%r\[0-9\]*\tatom.global.cas" 1 } } */
/* { dg-final { scan-assembler-times "shfl.idx.b32" 1 } } */
/* { dg-final { scan-assembler-times "vote.ballot.b32" 1 } } */