[gcc r11-10256] [nvptx] Add bar.warp.sync

[gcc r11-10256] [nvptx] Add bar.warp.sync

[Thomas Schwinge]

https://gcc.gnu.org/g:8b447b4be099ed5854666d53e6aa555b225403ac

commit r11-10256-g8b447b4be099ed5854666d53e6aa555b225403ac
Author: Tom de Vries <tdevries@suse.de>
Date:   Thu Jan 27 15:03:59 2022 +0100

    [nvptx] Add bar.warp.sync
    
    On a GT 1030 (sm_61), with driver version 470.94 I run into:
    ...
    FAIL: libgomp.oacc-c/../libgomp.oacc-c-c++-common/parallel-dims.c \
      -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 -foffload=nvptx-none \
      -O2 execution test
    ...
    which minimizes to the same test-case as listed in commit "[nvptx] Update
    default ptx isa to 6.3".
    
    The first divergent branch looks like:
    ...
      {
        .reg .u32 %x;
        mov.u32 %x,%tid.x;
        setp.ne.u32 %r59,%x,0;
      }
      @ %r59 bra $L15;
      mov.u64 %r48,%ar0;
      mov.u32 %r22,2;
      ld.u64 %r53,[%r48];
      mov.u32 %r55,%r22;
      mov.u32 %r54,1;
     $L15:
    ...
    and when inspecting the generated SASS, the branch is not setup as a divergent
    branch, but instead as a regular branch.
    
    This causes us to execute a shfl.sync insn in divergent mode, which is likely
    to cause trouble given a remark in the ptx isa version 6.3, which mentions
    that for .target sm_6x or below, all threads must excute the same
    shfl.sync instruction in convergence.
    
    Fix this by placing a "bar.warp.sync 0xffffffff" at the desired convergence
    point (in the example above, after $L15).
    
    Tested on x86_64 with nvptx accelerator.
    
    gcc/ChangeLog:
    
    2022-01-31  Tom de Vries  <tdevries@suse.de>
    
            * config/nvptx/nvptx.c (nvptx_single): Use nvptx_warpsync.
            * config/nvptx/nvptx.md (define_c_enum "unspecv"): Add
            UNSPECV_WARPSYNC.
            (define_insn "nvptx_warpsync"): New define_insn.
    
    (cherry picked from commit bba61d403d05202deb698b352a4faef3feb1f04d)

Diff:
---
 gcc/config/nvptx/nvptx.c  | 7 +++++++
 gcc/config/nvptx/nvptx.md | 7 +++++++
 2 files changed, 14 insertions(+)

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 7a7a9130e84..06861f6f405 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4503,6 +4503,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
   rtx_insn *neuter_start = NULL;
   rtx_insn *worker_label = NULL, *vector_label = NULL;
   rtx_insn *worker_jump = NULL, *vector_jump = NULL;
+  rtx_insn *warp_sync = NULL;
   for (mode = GOMP_DIM_WORKER; mode <= GOMP_DIM_VECTOR; mode++)
     if (GOMP_DIM_MASK (mode) & skip_mask)
       {
@@ -4535,11 +4536,15 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
 	if (tail_branch)
 	  {
 	    label_insn = emit_label_before (label, before);
+	    if (TARGET_PTX_6_0 && mode == GOMP_DIM_VECTOR)
+	      warp_sync = emit_insn_after (gen_nvptx_warpsync (), label_insn);
 	    before = label_insn;
 	  }
 	else
 	  {
 	    label_insn = emit_label_after (label, tail);
+	    if (TARGET_PTX_6_0 && mode == GOMP_DIM_VECTOR)
+	      warp_sync = emit_insn_after (gen_nvptx_warpsync (), label_insn);
 	    if ((mode == GOMP_DIM_VECTOR || mode == GOMP_DIM_WORKER)
 		&& CALL_P (tail) && find_reg_note (tail, REG_NORETURN, NULL))
 	      emit_insn_after (gen_exit (), label_insn);
@@ -4607,6 +4612,8 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
 		 setp.ne.u32 %rcond,%rcondu32,0;
 	  */
 	  rtx_insn *label = PREV_INSN (tail);
+	  if (label == warp_sync)
+	    label = PREV_INSN (label);
 	  gcc_assert (label && LABEL_P (label));
 	  rtx tmp = gen_reg_rtx (BImode);
 	  emit_insn_before (gen_movbi (tmp, const0_rtx),
diff --git a/gcc/config/nvptx/nvptx.md b/gcc/config/nvptx/nvptx.md
index 0f15609ee4b..c33b5d3279f 100644
--- a/gcc/config/nvptx/nvptx.md
+++ b/gcc/config/nvptx/nvptx.md
@@ -55,6 +55,7 @@
    UNSPECV_CAS
    UNSPECV_XCHG
    UNSPECV_BARSYNC
+   UNSPECV_WARPSYNC
    UNSPECV_MEMBAR
    UNSPECV_MEMBAR_CTA
    UNSPECV_DIM_POS
@@ -1718,6 +1719,12 @@
   }
   [(set_attr "predicable" "false")])
 
+(define_insn "nvptx_warpsync"
+  [(unspec_volatile [(const_int 0)] UNSPECV_WARPSYNC)]
+  "TARGET_PTX_6_0"
+  "\\tbar.warp.sync\\t0xffffffff;"
+  [(set_attr "predicable" "false")])
+
 (define_expand "memory_barrier"
   [(set (match_dup 0)
 	(unspec_volatile:BLK [(match_dup 0)] UNSPECV_MEMBAR))]