[og7] Update nvptx_fork/join barrier placement

[Cesar Philippidis <cesar@codesourcery.com>]

[-- Attachment #1: Type: text/plain, Size: 2994 bytes --]

Nvidia Volta GPUs now support warp-level synchronization. As such, the
semantics of legacy bar.sync instructions have slightly changed on newer
GPUs. The PTX JIT will now, occasionally, emit a warpsync instruction
immediately before a bar.sync for Volta GPUs. That implies that warps
must be convergent on entry to those threads barriers.

The problem in og7, and trunk, is that GCC emits barrier instructions at
the wrong spots. E.g., consider the following OpenACC parallel region:

  #pragma acc parallel loop worker
  for (i = 0; i < 10; i++)
    a[i] = i;

At -O2, GCC generates the following PTX code:

        {
                .reg.u32        %y;
                mov.u32 %y, %tid.y;
                setp.ne.u32     %r76, %y, 0;
        }
        {
                .reg.u32        %x;
                mov.u32 %x, %tid.x;
                setp.ne.u32     %r75, %x, 0;
        }
        @%r76   bra.uni $L6;
        @%r75   bra     $L7;
                mov.u64 %r67, %ar0;
        // fork 2;
                cvta.shared.u64 %r74, __oacc_bcast;
                st.u64  [%r74], %r67;
$L7:
$L6:
        @%r75   bra     $L5;
        // forked 2;
                bar.sync        0;
                cvta.shared.u64 %r73, __oacc_bcast;
                ld.u64  %r67, [%r73];
                mov.u32 %r62, %ntid.y;
                mov.u32 %r63, %tid.y;
                setp.gt.s32     %r68, %r63, 9;
        @%r68   bra     $L2;
                mov.u32 %r55, %r63;
                cvt.s64.s32     %r69, %r62;
                shl.b64 %r59, %r69, 2;
                cvt.s64.s32     %r70, %r55;
                shl.b64 %r71, %r70, 2;
                add.u64 %r58, %r67, %r71;
$L3:
                st.u32  [%r58], %r55;
                add.u32 %r55, %r55, %r62;
                add.u64 %r58, %r58, %r59;
                setp.le.s32     %r72, %r55, 9;
        @%r72   bra     $L3;
$L2:
                bar.sync        1;
        // joining 2;
$L5:
        // join 2;
        ret;

Note the bar.sync instructions placed immediately after the forked
comment and before the joining comment. The problem here is that branch
above the forked comment guarantees that the warps are not synchronous
(when vector_length > 1, which is always the case). Likewise, bar.sync
instruction before joining should be placed after label L5 in order to
allow all of the threads in the warp to reach it.

The attached patch teaches the nvptx to make those adjustments. It
doesn't cause any regressions on legacy GPUs, but it does resolve quite
a few failures with Volta in the libgomp execution tests. Therefore,
this patch doesn't include any new test cases. Part of this patch came
from my vector_length patch set that I posted last week. However, that
patch set didn't consider the placement of the joining barrier.

I've applied this patch to openacc-gcc-7-branch.

Tom, is a similar patch OK for trunk? The major difference between trunk
and og7 is that og7 changed the name of nvptx_warp_sync to nvptx_cta_sync.

Cesar

[-- Attachment #2: og7-barriers.diff --]
[-- Type: text/x-patch, Size: 1953 bytes --]

2018-03-08  Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* config/nvptx/nvptx.c (nvptx_single): Adjust placement of nvptx_fork
	and nvptx_join nutering labels.
	(nvptx_process_pars): Place the CTA barrier at the beginning of the
	join block.


diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index b16cf59575c..efc6161a6b0 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4056,6 +4056,15 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
 	return;
     }
 
+  /* NVPTX_BARSYNC barriers are placed immediately before NVPTX_JOIN
+     in order to ensure that all of the threads in a CTA reach the
+     barrier.  Don't nueter BLOCK if head is NVPTX_BARSYNC and tail is
+     NVPTX_JOIN.  */
+  if (from == to
+      && recog_memoized (head) == CODE_FOR_nvptx_barsync
+      && recog_memoized (tail) == CODE_FOR_nvptx_join)
+    return;
+
   /* Insert the vector test inside the worker test.  */
   unsigned mode;
   rtx_insn *before = tail;
@@ -4103,7 +4112,17 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
 	  br = gen_br_true (pred, label);
 	else
 	  br = gen_br_true_uni (pred, label);
-	emit_insn_before (br, head);
+
+	if (recog_memoized (head) == CODE_FOR_nvptx_forked
+	    && recog_memoized (NEXT_INSN (head)) == CODE_FOR_nvptx_barsync)
+	  {
+	    head = NEXT_INSN (head);
+	    emit_insn_after (br, head);
+	  }
+	else if (recog_memoized (head) == CODE_FOR_nvptx_barsync)
+	  emit_insn_after (br, head);
+	else
+	  emit_insn_before (br, head);
 
 	LABEL_NUSES (label)++;
 	if (tail_branch)
@@ -4325,7 +4344,7 @@ nvptx_process_pars (parallel *par)
 	{
 	  /* Insert begin and end synchronizations.  */
 	  emit_insn_after (nvptx_cta_sync (false), par->forked_insn);
-	  emit_insn_before (nvptx_cta_sync (true), par->joining_insn);
+	  emit_insn_before (nvptx_cta_sync (true), par->join_insn);
 	}
     }
   else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR))