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

diff --git a/gcc/config/nvptx/nvptx.cc b/gcc/config/nvptx/nvptx.cc
index afbad5b..4942f11 100644 (file)
--- a/gcc/config/nvptx/nvptx.cc
+++ b/gcc/config/nvptx/nvptx.cc
@@ -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);

diff --git a/gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c b/gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c
index 1bc0ada..77cffc4 100644 (file)
--- a/gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c
+++ b/gcc/testsuite/gcc.target/nvptx/uniform-simt-1.c
@@ -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 } } */

diff --git a/gcc/testsuite/gcc.target/nvptx/uniform-simt-2.c b/gcc/testsuite/gcc.target/nvptx/uniform-simt-2.c
new file mode 100644 (file)
index 0000000..0f1e4e7
--- /dev/null
+++ b/gcc/testsuite/gcc.target/nvptx/uniform-simt-2.c
@@ -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 } } */