[nvptx] Disable warp sync in simt region

I ran into a hang for this code:
...
  #pragma omp target map(tofrom: counter_N0)
  #pragma omp simd
  for (int i = 0 ; i < 1 ; i++ )
    {
      #pragma omp atomic update
      counter_N0 = counter_N0 + 1 ;
    }
...

This has to do with the nature of -muniform-simt.  It has two modes of
operation: inside and outside an SIMT region.

Outside an SIMT region, a warp pretends to execute a single thread, but
actually executes in all threads, to keep the local registers in all threads
consistent.  This approach works unless the insn that is executed is a syscall
or an atomic insn.  In that case, the insn is predicated, such that it
executes in only one thread.  If the predicated insn writes a result to a
register, then that register is propagated to the other threads, after which
the local registers in all threads are consistent again.

Inside an SIMT region, a warp executes in all threads.  However, the
predication and propagation for syscalls and atomic insns is also present
here, because nvptx_reorg_uniform_simt works on all code.  Care has been taken
though to ensure that the predication and propagation is a nop.  That is,
inside an SIMT region:
- the predicate evalutes to true for each thread, and
- the propagation insn copies a register from each thread to the same thread.

That works fine, until we use -mptx=6.0, and instead of using the deprecated
warp propagation insn shfl, we start using shfl.sync:
...
  @%r33 atom.add.u32		_, [%r29], 1;
	shfl.sync.idx.b32	%r30, %r30, %r32, 31, 0xffffffff;
...

The shfl.sync specifies a member mask indicating all threads, but given that
the loop only has a single iteration, only thread 0 will execute the insn,
where it will hang waiting for the other threads.

Fix this by predicating the shfl.sync (and likewise, bar.warp.sync and the
uniform warp check) such that it only executes outside the SIMT region.

Tested on x86_64 with nvptx accelerator.

gcc/ChangeLog:

2022-03-08  Tom de Vries  <tdevries@suse.de>

	PR target/104783
	* config/nvptx/nvptx.cc (nvptx_init_unisimt_predicate)
	(nvptx_output_unisimt_switch): Handle unisimt_outside_simt_predicate.
	(nvptx_get_unisimt_outside_simt_predicate): New function.
	(predicate_insn): New function, factored out of ...
	(nvptx_reorg_uniform_simt): ... here.  Predicate all emitted insns.
	* config/nvptx/nvptx.h (struct machine_function): Add
	unisimt_outside_simt_predicate field.
	* config/nvptx/nvptx.md (define_insn "nvptx_warpsync")
	(define_insn "nvptx_uniform_warp_check"): Make predicable.

libgomp/ChangeLog:

2022-03-10  Tom de Vries  <tdevries@suse.de>

	* testsuite/libgomp.c/pr104783.c: New test.

gcc/config/nvptx/nvptx.cc

@@ -1364,6 +1364,13 @@ nvptx_init_unisimt_predicate (FILE *file)
       int master = REGNO (cfun->machine->unisimt_master);
       int pred = REGNO (cfun->machine->unisimt_predicate);
       fprintf (file, "\t\tld.shared.u32 %%r%d, [%%r%d];\n", master, loc);
+      if (cfun->machine->unisimt_outside_simt_predicate)
+	{
+	  int pred_outside_simt
+	    = REGNO (cfun->machine->unisimt_outside_simt_predicate);
+	  fprintf (file, "\t\tsetp.eq.u32 %%r%d, %%r%d, 0;\n",
+		   pred_outside_simt, master);
+	}
       fprintf (file, "\t\tmov.u32 %%ustmp0, %%laneid;\n");
       /* Compute 'master lane index' as 'laneid & __nvptx_uni[tid.y]'.  */
       fprintf (file, "\t\tand.b32 %%r%d, %%r%d, %%ustmp0;\n", master, master);
@@ -1589,6 +1596,13 @@ nvptx_output_unisimt_switch (FILE *file, bool entering)
   fprintf (file, "\t{\n");
   fprintf (file, "\t\t.reg.u32 %%ustmp2;\n");
   fprintf (file, "\t\tmov.u32 %%ustmp2, %d;\n", entering ? -1 : 0);
+  if (cfun->machine->unisimt_outside_simt_predicate)
+    {
+      int pred_outside_simt
+	= REGNO (cfun->machine->unisimt_outside_simt_predicate);
+      fprintf (file, "\t\tmov.pred %%r%d, %d;\n", pred_outside_simt,
+	       entering ? 0 : 1);
+    }
   if (!crtl->is_leaf)
     {
       int loc = REGNO (cfun->machine->unisimt_location);
@@ -3242,6 +3256,13 @@ nvptx_get_unisimt_predicate ()
   return pred ? pred : pred = gen_reg_rtx (BImode);
 }
 
+static rtx
+nvptx_get_unisimt_outside_simt_predicate ()
+{
+  rtx &pred = cfun->machine->unisimt_outside_simt_predicate;
+  return pred ? pred : pred = gen_reg_rtx (BImode);
+}
+
 /* Return true if given call insn references one of the functions provided by
    the CUDA runtime: malloc, free, vprintf.  */
 
@@ -3286,6 +3307,16 @@ nvptx_unisimt_handle_set (rtx set, rtx_insn *insn, rtx master)
   return false;
 }
 
+static void
+predicate_insn (rtx_insn *insn, rtx pred)
+{
+  rtx pat = PATTERN (insn);
+  pred = gen_rtx_NE (BImode, pred, const0_rtx);
+  pat = gen_rtx_COND_EXEC (VOIDmode, pred, pat);
+  bool changed_p = validate_change (insn, &PATTERN (insn), pat, false);
+  gcc_assert (changed_p);
+}
+
 /* Adjust code for uniform-simt code generation variant by making atomics and
    "syscalls" conditionally executed, and inserting shuffle-based propagation
    for registers being set.  */
@@ -3352,10 +3383,16 @@ nvptx_reorg_uniform_simt ()
 	}
 
       rtx pred = nvptx_get_unisimt_predicate ();
-      pred = gen_rtx_NE (BImode, pred, const0_rtx);
-      pat = gen_rtx_COND_EXEC (VOIDmode, pred, pat);
-      bool changed_p = validate_change (insn, &PATTERN (insn), pat, false);
-      gcc_assert (changed_p);
+      predicate_insn (insn, pred);
+
+      pred = NULL_RTX;
+      for (rtx_insn *post = NEXT_INSN (insn); post != next;
+	   post = NEXT_INSN (post))
+	{
+	  if (pred == NULL_RTX)
+	    pred = nvptx_get_unisimt_outside_simt_predicate ();
+	  predicate_insn (post, pred);
+	}
     }
 }
 

gcc/config/nvptx/nvptx.h

@@ -226,6 +226,7 @@ struct GTY(()) machine_function
   rtx sync_bar; /* Synchronization barrier ID for vectors.  */
   rtx unisimt_master; /* 'Master lane index' for -muniform-simt.  */
   rtx unisimt_predicate; /* Predicate for -muniform-simt.  */
+  rtx unisimt_outside_simt_predicate; /* Predicate for -muniform-simt.  */
   rtx unisimt_location; /* Mask location for -muniform-simt.  */
   /* The following two fields hold the maximum size resp. alignment required
      for per-lane storage in OpenMP SIMD regions.  */

gcc/config/nvptx/nvptx.md

@@ -2268,25 +2268,28 @@
 (define_insn "nvptx_warpsync"
   [(unspec_volatile [(const_int 0)] UNSPECV_WARPSYNC)]
   "TARGET_PTX_6_0"
-  "\\tbar.warp.sync\\t0xffffffff;"
-  [(set_attr "predicable" "false")])
+  "%.\\tbar.warp.sync\\t0xffffffff;")
 
 (define_insn "nvptx_uniform_warp_check"
   [(unspec_volatile [(const_int 0)] UNSPECV_UNIFORM_WARP_CHECK)]
   ""
   {
-    output_asm_insn ("{", NULL);
-    output_asm_insn ("\\t"	 ".reg.b32"	   "\\t" "act;", NULL);
-    output_asm_insn ("\\t"	 "vote.ballot.b32" "\\t" "act,1;", NULL);
-    output_asm_insn ("\\t"	 ".reg.pred"	   "\\t" "uni;", NULL);
-    output_asm_insn ("\\t"	 "setp.eq.b32"	   "\\t" "uni,act,0xffffffff;",
-		     NULL);
-    output_asm_insn ("@ !uni\\t" "trap;", NULL);
-    output_asm_insn ("@ !uni\\t" "exit;", NULL);
-    output_asm_insn ("}", NULL);
+    const char *insns[] = {
+      "{",
+      "\\t"	      ".reg.b32"	"\\t" "act;",
+      "%.\\t"	      "vote.ballot.b32" "\\t" "act,1;",
+      "\\t"	      ".reg.pred"	"\\t" "do_abort;",
+      "\\t"	      "mov.pred"	"\\t" "do_abort,0;",
+      "%.\\t"	      "setp.ne.b32"	"\\t" "do_abort,act,0xffffffff;",
+      "@ do_abort\\t" "trap;",
+      "@ do_abort\\t" "exit;",
+      "}",
+      NULL
+    };
+    for (const char **p = &insns[0]; *p != NULL; p++)
+      output_asm_insn (*p, NULL);
     return "";
-  }
-  [(set_attr "predicable" "false")])
+  })
 
 (define_expand "memory_barrier"
   [(set (match_dup 0)

libgomp/testsuite/libgomp.c/pr104783.c

@@ -0,0 +1,18 @@
+int
+main (void)
+{
+  unsigned val = 0;
+
+#pragma omp target map(tofrom: val)
+#pragma omp simd
+  for (int i = 0 ; i < 1 ; i++)
+    {
+#pragma omp atomic update
+      val = val + 1;
+    }
+
+  if (val != 1)
+    __builtin_abort ();
+
+  return 0;
+}