Re: [og7] Update nvptx_fork/join barrier placement

[Cesar Philippidis]

On 03/19/2018 10:02 AM, Tom de Vries wrote:
> On 03/19/2018 03:55 PM, Cesar Philippidis wrote:
>>> Note that this changes ordering of the vector-neutering jump and
>>> worker-neutering jump at the end. In principle, this should not be
>>> harmful, but it violates the invariant that vector-neutering
>>> branch-around code should be as short-lived as possible. So, this needs
>>> to be fixed.
>>>
>>> I've found this issue by adding verification of the neutering, as
>>> attached below.
>> ACK, thanks. I'll take a closer look at this.
> 
> I've got a tentative patch at
> https://gcc.gnu.org/bugzilla/attachment.cgi?id=43707 ( PR84952 -
> "[nvptx] bar.sync generated in divergent code" ).

I attached my WIP patch. But, given that you've spent a lot of time on
this, I'll let you continue working on it. Just remember to backport any
fix to og7.

Thanks,
Cesar

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index a6f444340fd..0d288cb81ba 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4037,6 +4037,22 @@ 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;
+
+  /* Adjust HEAD to point to the NVPTX_JOIN instruction after a
+     NVPTX_BARSYNC, so that any successive state neutering code does
+     not get placed before the dummy JOIN comment. */
+  if (recog_memoized (head) == CODE_FOR_nvptx_barsync
+      && recog_memoized (NEXT_INSN (head)) == CODE_FOR_nvptx_join)
+    head = NEXT_INSN (head);
+
   /* Insert the vector test inside the worker test.  */
   unsigned mode;
   rtx_insn *before = tail;
@@ -4057,7 +4073,23 @@ 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_join)
+	  {
+	    if (recog_memoized (NEXT_INSN (head)) == CODE_FOR_br_true_uni
+		&& mode == GOMP_DIM_VECTOR)
+	      emit_insn_after (br, NEXT_INSN (head));
+	    else
+	      emit_insn_after (br, head);
+	  }
+	else
+	  emit_insn_before (br, head);
 
 	LABEL_NUSES (label)++;
 	if (tail_branch)
@@ -4276,7 +4308,7 @@ nvptx_process_pars (parallel *par)
       nvptx_wpropagate (true, par->forked_block, par->fork_insn);
       /* Insert begin and end synchronizations.  */
       emit_insn_after (nvptx_wsync (false), par->forked_insn);
-      emit_insn_before (nvptx_wsync (true), par->joining_insn);
+      emit_insn_before (nvptx_wsync (true), par->join_insn);
     }
   else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR))
     nvptx_vpropagate (par->forked_block, par->forked_insn);