@@ -3989,13 +3989,14 @@ nvptx_shared_propagate (bool pre_p, bool is_call, basic_block block,
return empty;
}
-/* Emit a CTA-level synchronization barrier. We use different
- markers for before and after synchronizations. */
+/* Emit a CTA-level synchronization barrier (bar.sync). LOCK is the
+ barrier number, which is an integer or a register. THREADS is the
+ number of threads controlled by the barrier. */
static rtx
-nvptx_cta_sync (bool after)
+nvptx_cta_sync (rtx lock, int threads)
{
- return gen_nvptx_barsync (GEN_INT (after));
+ return gen_nvptx_barsync (lock, GEN_INT (threads));
}
#if WORKAROUND_PTXJIT_BUG
@@ -4355,6 +4356,8 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
/* Includes worker mode, do spill & fill. By construction
we should never have worker mode only. */
broadcast_data_t data;
+ rtx barrier = GEN_INT (0);
+ int threads = 0;
data.base = oacc_bcast_sym;
data.ptr = 0;
@@ -4367,14 +4370,14 @@ nvptx_single (unsigned mask, basic_block from, basic_block to)
false),
before);
/* Barrier so other workers can see the write. */
- emit_insn_before (nvptx_cta_sync (false), tail);
+ emit_insn_before (nvptx_cta_sync (barrier, threads), tail);
data.offset = 0;
emit_insn_before (nvptx_gen_shared_bcast (pvar, PM_write, 0, &data,
false), tail);
/* 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_cta_sync (false), tail);
+ emit_insn_before (nvptx_cta_sync (barrier, threads), tail);
}
extract_insn (tail);
@@ -4496,12 +4499,15 @@ nvptx_process_pars (parallel *par)
bool empty = nvptx_shared_propagate (true, is_call,
par->forked_block, par->fork_insn,
false);
+ rtx barrier = GEN_INT (0);
+ int threads = 0;
if (!empty || !is_call)
{
/* Insert begin and end synchronizations. */
- emit_insn_before (nvptx_cta_sync (false), par->forked_insn);
- emit_insn_before (nvptx_cta_sync (false), par->join_insn);
+ emit_insn_before (nvptx_cta_sync (barrier, threads),
+ par->forked_insn);
+ emit_insn_before (nvptx_cta_sync (barrier, threads), par->join_insn);
}
}
else if (par->mask & GOMP_DIM_MASK (GOMP_DIM_VECTOR))
@@ -1455,10 +1455,16 @@
[(set_attr "atomic" "true")])
(define_insn "nvptx_barsync"
- [(unspec_volatile [(match_operand:SI 0 "const_int_operand" "")]
+ [(unspec_volatile [(match_operand:SI 0 "nvptx_nonmemory_operand" "Ri")
+ (match_operand:SI 1 "const_int_operand")]
UNSPECV_BARSYNC)]
""
- "\\tbar.sync\\t%0;"
+ {
+ if (INTVAL (operands[1]) == 0)
+ return "\\tbar.sync\\t%0;";
+ else
+ return "\\tbar.sync\\t%0, %1;";
+ }
[(set_attr "predicable" "false")])
(define_expand "memory_barrier"
From: Tom de Vries <tom@codesourcery.com> This patch extends the nvptx_barsync pattern to accept an integeral parameter to specify how many threads utilize the requested barrier. This is necessary variable length vectors, as each large vector will require own thread barrier. 2018-XX-YY Tom de Vries <tdevries@suse.de> Cesar Philippidis <cesar@codesourcery.com> gcc/ * config/nvptx/nvptx.md (nvptx_barsync): Add and handle operand. * config/nvptx/nvptx.c (nvptx_cta_sync): Change arguments to take in a lock and thread count. Update call to gen_nvptx_barsync. (nvptx_single, nvptx_process_pars): Update calls to nvptx_cta_sync. (cherry picked from openacc-gcc-7-branch commit 49b69b4002496c9f5759d933e0a6233660434e70)