[nvptx] Fix bar.sync position

2018-03-20  Tom de Vries  <tom@codesourcery.com>

	PR target/84952
	* config/nvptx/nvptx.c (nvptx_single): Don't neuter bar.sync.
	(nvptx_process_pars): Emit bar.sync asap and alap.

From-SVN: r258676

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 460802d8b643027c7591c755a53268fa210b5fdb..c1ac2f9228a889cb7ec81de0dcca7598a1a4a4c0 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2018-03-20  Tom de Vries  <tom@codesourcery.com>
+
+       PR target/84952
+       * config/nvptx/nvptx.c (nvptx_single): Don't neuter bar.sync.
+       (nvptx_process_pars): Emit bar.sync asap and alap.
+
 2018-03-20  Tom de Vries  <tom@codesourcery.com>
 
        PR target/84954

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 7b0b182deb2283bdf1235ac3e7de7155d954cd24..1ba27e31ef088d3c9bdf3dc42c0a355e9ec53cb2 100644 (file)
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -3969,7 +3969,9 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
   while (true)
     {
       /* Find first insn of from block.  */
-      while (head != BB_END (from) && !INSN_P (head))
+      while (head != BB_END (from)
+            && (!INSN_P (head)
+                || recog_memoized (head) == CODE_FOR_nvptx_barsync))
        head = NEXT_INSN (head);
 
       if (from == to)
@@ -4018,6 +4020,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
        {
        default:
          break;
+       case CODE_FOR_nvptx_barsync:
        case CODE_FOR_nvptx_fork:
        case CODE_FOR_nvptx_forked:
        case CODE_FOR_nvptx_joining:
@@ -4275,8 +4278,8 @@ nvptx_process_pars (parallel *par)
       nvptx_wpropagate (false, par->forked_block, par->forked_insn);
       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 (false), par->forked_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);