[nvptx] Only use one logical barrier resource

For openacc loops, we generate this style of code:
...
        @%r41   bra.uni $L5;
        @%r40   bra     $L6;
                mov.u64 %r32, %ar0;
                cvta.shared.u64 %r39, __worker_bcast;
                st.u64  [%r39], %r32;
$L6:
$L5:
                bar.sync        0;
        @%r40   bra     $L4;
                cvta.shared.u64 %r38, __worker_bcast;
                ld.u64  %r32, [%r38];
                ...
$L4:
                bar.sync        1;
...

The first barrier is there to ensure that no thread reads the broadcast buffer
before it's written.  The second barrier is there to ensure that no thread
overwrites the broadcast buffer before all threads have read it (as well as
implementing the obligatory synchronization after a worker loop).

We've been using the logical barrier resources '0' and '1' for these two
barriers, but there's no reason why we can't use the same one.

Use logical barrier resource '0' for both barriers, making the openacc
implementation claim less resources.

Build and reg-tested on x86_64 with nvptx accelerator.

2018-12-19  Tom de Vries  <tdevries@suse.de>

	* config/nvptx/nvptx.c (nvptx_single): Always pass false to
	nvptx_wsync.
	(nvptx_process_pars): Likewise.

From-SVN: r267257

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index 82f381a0b5a5d8544f415e2eb307c5e66658a9f4..9596b26860c925ad99f8b0f9bf8a9cbe6b49b207 100644 (file)
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,3 +1,9 @@
+2018-12-19  Tom de Vries  <tdevries@suse.de>
+
+       * config/nvptx/nvptx.c (nvptx_single): Always pass false to
+       nvptx_wsync.
+       (nvptx_process_pars): Likewise.
+
 2018-12-19  Tom de Vries  <tdevries@suse.de>
 
        * config/nvptx/nvptx.c (nvptx_previous_fndecl): Declare.

diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 9f834d35200d43b763e281e5a4183a1b84baff6a..a354811194c57064bbe0e67f6cb6cda0cec96aad 100644 (file)
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -4351,7 +4351,7 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
          /* This barrier is needed to avoid worker zero clobbering
             the broadcast buffer before all the other workers have
             had a chance to read this instance of it.  */
-         emit_insn_before (nvptx_wsync (true), tail);
+         emit_insn_before (nvptx_wsync (false), tail);
        }
 
       extract_insn (tail);
@@ -4476,7 +4476,7 @@ nvptx_process_pars (parallel *par)
        {
          /* Insert begin and end synchronizations.  */
          emit_insn_before (nvptx_wsync (false), par->forked_insn);
-         emit_insn_before (nvptx_wsync (true), par->join_insn);
+         emit_insn_before (nvptx_wsync (false), par->join_insn);
        }
     }
   else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR))