public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] openacc: Shared memory layout optimisation
@ 2021-05-13 16:15 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:15 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:01c5168502e2ee0891f7316047f7f088827770fa
commit 01c5168502e2ee0891f7316047f7f088827770fa
Author: Julian Brown <julian@codesourcery.com>
Date: Mon Jan 27 06:14:05 2020 -0800
openacc: Shared memory layout optimisation
This patch implements an algorithm to lay out local data-share (LDS)
space. It currently works for AMD GCN. At the moment, LDS is used for
three things:
1. Gang-private variables
2. Reduction temporaries (accumulators)
3. Broadcasting for worker partitioning
After the patch is applied, (2) and (3) are placed at preallocated
locations in LDS, and (1) continues to be handled by the backend (as it
is at present prior to this patch being applied). LDS now looks like this:
+--------------+ (gang local size + 1024, = 1536)
| free space |
| ... |
| - - - - - - -|
| worker bcast |
+--------------+
| reductions |
+--------------+ <<< -mgang-local-size=<number> (def. 512)
| gang private |
| vars |
+--------------+ (32)
| low LDS vars |
+--------------+ LDS base
So, gang-private space is fixed at a constant amount at compile time
(which can be increased with a command-line switch if necessary
for some given code). The layout algorithm takes out a slice of the
remainder of usable space for reduction vars, and uses the rest for
worker partitioning.
The partitioning algorithm works as follows.
1. An "adjacency" set is built up for each basic block that might
do a broadcast. This is calculated by starting at each such block,
and doing a recursive DFS walk over successors to find the next
block (or blocks) that *also* does a broadcast
(dfs_broadcast_reachable_1).
2. The adjacency set is inverted to get adjacent predecessor blocks also.
3. Blocks that will perform a broadcast are sorted by size of that
broadcast: the biggest blocks are handled first.
4. A splay tree structure is used to calculate the spans of LDS memory
that are already allocated by the blocks adjacent to this one
(merge_ranges{,_1}.
5. The current block's broadcast space is allocated from the first free
span not allocated in the splay tree structure calculated above
(first_fit_range). This seems to work quite nicely and efficiently
with the splay tree structure.
6. Continue with the next-biggest broadcast block until we're done.
In this way, "adjacent" broadcasts will not use the same piece of
LDS memory.
2020-07-15 Julian Brown <julian@codesourcery.com>
gcc/
* config/gcn/gcn-protos.h (gcn_goacc_adjust_private_decl): Update
prototype.
* config/gcn/gcn-tree.c (gcn_goacc_get_worker_red_decl): Use
preallocated block of LDS memory.
(gcn_goacc_create_propagation_record): Add OFFSET parameter, and return
temporary LDS space at that offset. Return pointer in "sender" case.
(gcn_goacc_adjust_private_decl): Return var.
* config/gcn/gcn.c (acc_lds_size, gangprivate_hwm, lds_allocs): New
global vars.
(ACC_LDS_SIZE): Define as acc_lds_size.
(gcn_init_machine_status): Don't initialise lds_allocated and
lds_allocs fields of machine function struct.
(gcn_option_override): Handle default size for gang-private variables
and -mgang-local-size option.
(gcn_expand_prologue): Use LDS_SIZE instead of LDS_SIZE-1 when
initialising M0_REG.
(gcn_shared_mem_layout): New function.
(gcn_print_lds_decl): Update comment. Use global lds_allocs map and
gangprivate_hwm variable.
(TARGET_GOACC_SHARED_MEM_LAYOUT): Define target hook.
* config/gcn/gcn.h (machine_function): Remove lds_allocated,
lds_allocs. Add reduction_base, reduction_limit.
* config/gcn/gcn.opt (gang_local_size_opt): New global.
(mgang-local-size=): New option.
* config/nvptx/nvptx.c (nvptx_goacc_adjust_private_decl): Return var.
* doc/tm.texi.in (TARGET_GOACC_SHARED_MEM_LAYOUT): Place documentation
hook.
* doc/tm.texi: Regenerate.
* omp-offload.c (addr_expr_rewrite_info): Change adjusted_vars to a
hash_map.
(rewrite_addr_expr): Rewrite VAR_DECLs also.
(default_goacc_create_propagation_record): Add OFFSET parameter.
(execute_oacc_gimple_workers): Calculate per-function reduction
temporary and private-variable size. Call OpenACC shared_mem_layout
hook. Move num_workers==1 handling here.
(execute_oacc_device_lower): Fix for adjusted_vars being a hash_map
rather than a hash_set.
(pass_oacc_gimple_workers::gate): Remove num_workers==1 handling from
here. Enable pass for all OpenACC routines in order to call shared
memory-layout hook.
* omp-sese.c (targhooks.h, diagnostic-core.h): Add includes.
(build_sender_ref): Handle sender_decl being pointer.
(worker_single_copy): Add PLACEMENT and ISOLATE_BROADCASTS parameters.
Pass placement argument to create_propagation_record hook invocations.
Handle sender_decl being pointer and isolate_broadcasts inserting extra
barriers.
(blk_offset_map_t): Add typedef.
(neuter_worker_single): Add BLK_OFFSET_MAP parameter. Pass
preallocated range to worker_single_copy call.
(dfs_broadcast_reachable_1): New function.
(idx_decl_pair_t, used_range_vec_t): New typedefs.
(sort_size_descending): New function.
(addr_range): New class.
(splay_tree_compare_addr_range, splay_tree_free_key, first_fit_range,
merge_ranges_1, merge_ranges): New functions.
(oacc_do_neutering): Add BOUNDS_LO, BOUNDS_HI parameters. Arrange
layout of shared memory for broadcast operations.
* omp-sese.h (oacc_do_neutering): Update prototype.
* target.def (adjust_private_decl): Change return type to tree.
(create_propagation_record): Add OFFSET parameter.
(shared_mem_layout): New hook.
* targhooks.h (default_goacc_create_propagation_record): Update
prototype.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/broadcast-many.c: New test.
Diff:
---
gcc/ChangeLog.omp | 66 +++++
gcc/config/gcn/gcn-protos.h | 5 +-
gcc/config/gcn/gcn-tree.c | 48 +--
gcc/config/gcn/gcn.c | 82 ++++--
gcc/config/gcn/gcn.h | 5 +-
gcc/config/gcn/gcn.opt | 6 +
gcc/config/nvptx/nvptx.c | 6 +-
gcc/doc/tm.texi | 11 +-
gcc/doc/tm.texi.in | 2 +
gcc/omp-offload.c | 152 ++++++++--
gcc/omp-sese.c | 326 ++++++++++++++++++++-
gcc/omp-sese.h | 2 +-
gcc/target.def | 14 +-
gcc/targhooks.h | 3 +-
libgomp/ChangeLog.omp | 4 +
.../libgomp.oacc-c-c++-common/broadcast-many.c | 79 +++++
16 files changed, 719 insertions(+), 92 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 500a58e5fd3..4f680f0ddd2 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,69 @@
+2020-07-15 Julian Brown <julian@codesourcery.com>
+
+ * config/gcn/gcn-protos.h (gcn_goacc_adjust_private_decl): Update
+ prototype.
+ * config/gcn/gcn-tree.c (gcn_goacc_get_worker_red_decl): Use
+ preallocated block of LDS memory.
+ (gcn_goacc_create_propagation_record): Add OFFSET parameter, and return
+ temporary LDS space at that offset. Return pointer in "sender" case.
+ (gcn_goacc_adjust_private_decl): Return var.
+ * config/gcn/gcn.c (acc_lds_size, gangprivate_hwm, lds_allocs): New
+ global vars.
+ (ACC_LDS_SIZE): Define as acc_lds_size.
+ (gcn_init_machine_status): Don't initialise lds_allocated and
+ lds_allocs fields of machine function struct.
+ (gcn_option_override): Handle default size for gang-private variables
+ and -mgang-local-size option.
+ (gcn_expand_prologue): Use LDS_SIZE instead of LDS_SIZE-1 when
+ initialising M0_REG.
+ (gcn_shared_mem_layout): New function.
+ (gcn_print_lds_decl): Update comment. Use global lds_allocs map and
+ gangprivate_hwm variable.
+ (TARGET_GOACC_SHARED_MEM_LAYOUT): Define target hook.
+ * config/gcn/gcn.h (machine_function): Remove lds_allocated,
+ lds_allocs. Add reduction_base, reduction_limit.
+ * config/gcn/gcn.opt (gang_local_size_opt): New global.
+ (mgang-local-size=): New option.
+ * config/nvptx/nvptx.c (nvptx_goacc_adjust_private_decl): Return var.
+ * doc/tm.texi.in (TARGET_GOACC_SHARED_MEM_LAYOUT): Place documentation
+ hook.
+ * doc/tm.texi: Regenerate.
+ * omp-offload.c (addr_expr_rewrite_info): Change adjusted_vars to a
+ hash_map.
+ (rewrite_addr_expr): Rewrite VAR_DECLs also.
+ (default_goacc_create_propagation_record): Add OFFSET parameter.
+ (execute_oacc_gimple_workers): Calculate per-function reduction
+ temporary and private-variable size. Call OpenACC shared_mem_layout
+ hook. Move num_workers==1 handling here.
+ (execute_oacc_device_lower): Fix for adjusted_vars being a hash_map
+ rather than a hash_set.
+ (pass_oacc_gimple_workers::gate): Remove num_workers==1 handling from
+ here. Enable pass for all OpenACC routines in order to call shared
+ memory-layout hook.
+ * omp-sese.c (targhooks.h, diagnostic-core.h): Add includes.
+ (build_sender_ref): Handle sender_decl being pointer.
+ (worker_single_copy): Add PLACEMENT and ISOLATE_BROADCASTS parameters.
+ Pass placement argument to create_propagation_record hook invocations.
+ Handle sender_decl being pointer and isolate_broadcasts inserting extra
+ barriers.
+ (blk_offset_map_t): Add typedef.
+ (neuter_worker_single): Add BLK_OFFSET_MAP parameter. Pass
+ preallocated range to worker_single_copy call.
+ (dfs_broadcast_reachable_1): New function.
+ (idx_decl_pair_t, used_range_vec_t): New typedefs.
+ (sort_size_descending): New function.
+ (addr_range): New class.
+ (splay_tree_compare_addr_range, splay_tree_free_key, first_fit_range,
+ merge_ranges_1, merge_ranges): New functions.
+ (oacc_do_neutering): Add BOUNDS_LO, BOUNDS_HI parameters. Arrange
+ layout of shared memory for broadcast operations.
+ * omp-sese.h (oacc_do_neutering): Update prototype.
+ * target.def (adjust_private_decl): Change return type to tree.
+ (create_propagation_record): Add OFFSET parameter.
+ (shared_mem_layout): New hook.
+ * targhooks.h (default_goacc_create_propagation_record): Update
+ prototype.
+
2020-07-15 Julian Brown <julian@codesourcery.com>
* omp-offload.c (pass_oacc_gimple_workers::gate): Disable worker
diff --git a/gcc/config/gcn/gcn-protos.h b/gcc/config/gcn/gcn-protos.h
index 61a2b00e936..c1a31aa1b49 100644
--- a/gcc/config/gcn/gcn-protos.h
+++ b/gcc/config/gcn/gcn-protos.h
@@ -39,8 +39,9 @@ extern rtx gcn_full_exec_reg ();
extern rtx gcn_gen_undef (machine_mode);
extern bool gcn_global_address_p (rtx);
extern tree gcn_goacc_create_propagation_record (tree record_type, bool sender,
- const char *name);
-extern void gcn_goacc_adjust_private_decl (tree var, int level);
+ const char *name,
+ unsigned HOST_WIDE_INT offset);
+extern tree gcn_goacc_adjust_private_decl (tree var, int level);
extern void gcn_goacc_reduction (gcall *call);
extern bool gcn_hard_regno_rename_ok (unsigned int from_reg,
unsigned int to_reg);
diff --git a/gcc/config/gcn/gcn-tree.c b/gcc/config/gcn/gcn-tree.c
index 90be7f15151..fb30a8627f4 100644
--- a/gcc/config/gcn/gcn-tree.c
+++ b/gcc/config/gcn/gcn-tree.c
@@ -329,14 +329,12 @@ gcn_goacc_get_worker_red_decl (tree type, unsigned offset)
}
else
{
- char name[50];
- sprintf (name, ".oacc_reduction_%u", offset);
- tree decl = create_tmp_var_raw (var_type, name);
+ gcc_assert (offset
+ < (machfun->reduction_limit - machfun->reduction_base));
+ tree ptr_type = build_pointer_type (var_type);
+ tree addr = build_int_cst (ptr_type, machfun->reduction_base + offset);
- DECL_CONTEXT (decl) = NULL_TREE;
- TREE_STATIC (decl) = 1;
-
- varpool_node::finalize_decl (decl);
+ tree decl = build_simple_mem_ref (addr);
vec_safe_grow_cleared (machfun->reduc_decls, offset + 1, true);
(*machfun->reduc_decls)[offset] = decl;
@@ -555,34 +553,36 @@ gcn_goacc_reduction (gcall *call)
tree
gcn_goacc_create_propagation_record (tree record_type, bool sender,
- const char *name)
+ const char *name,
+ unsigned HOST_WIDE_INT offset)
{
- tree type = record_type;
-
- TYPE_ADDR_SPACE (type) = ADDR_SPACE_LDS;
+ tree type = build_qualified_type (record_type,
+ TYPE_QUALS_NO_ADDR_SPACE (record_type)
+ | ENCODE_QUAL_ADDR_SPACE (ADDR_SPACE_LDS));
if (!sender)
- type = build_pointer_type (type);
-
- tree decl = create_tmp_var_raw (type, name);
-
- if (sender)
{
- DECL_CONTEXT (decl) = NULL_TREE;
- TREE_STATIC (decl) = 1;
+ tree ptr_type = build_pointer_type (type);
+ return create_tmp_var_raw (ptr_type, name);
}
- if (sender)
- varpool_node::finalize_decl (decl);
+ if (record_type == char_type_node)
+ offset = 1;
+
+ gcc_assert (cfun);
- return decl;
+ machine_function *machfun = cfun->machine;
+ unsigned HOST_WIDE_INT size = tree_to_uhwi (TYPE_SIZE_UNIT (record_type));
+
+ tree ptr_type = build_pointer_type (type);
+ return build_int_cst (ptr_type, offset);
}
-void
+tree
gcn_goacc_adjust_private_decl (tree var, int level)
{
if (level != GOMP_DIM_GANG)
- return;
+ return var;
tree type = TREE_TYPE (var);
tree lds_type = build_qualified_type (type,
@@ -601,6 +601,8 @@ gcn_goacc_adjust_private_decl (tree var, int level)
if (machfun)
machfun->use_flat_addressing = true;
+
+ return var;
}
/* }}} */
diff --git a/gcc/config/gcn/gcn.c b/gcc/config/gcn/gcn.c
index 3741853e76f..017e24a2474 100644
--- a/gcc/config/gcn/gcn.c
+++ b/gcc/config/gcn/gcn.c
@@ -72,14 +72,21 @@ int gcn_isa = 3; /* Default to GCN3. */
We want to permit full occupancy, so size accordingly. */
+/* Use this as a default, but allow it to grow if the user requests a large
+ amount of gang-private shared-memory space. */
+static int acc_lds_size = 0x600;
+
#define OMP_LDS_SIZE 0x600 /* 0x600 is 1/40 total, rounded down. */
-#define ACC_LDS_SIZE 32768 /* Half of the total should be fine. */
+#define ACC_LDS_SIZE acc_lds_size
#define OTHER_LDS_SIZE 65536 /* If in doubt, reserve all of it. */
#define LDS_SIZE (flag_openacc ? ACC_LDS_SIZE \
: flag_openmp ? OMP_LDS_SIZE \
: OTHER_LDS_SIZE)
+static int gangprivate_hwm = 32;
+static hash_map<tree, int> lds_allocs;
+
/* The number of registers usable by normal non-kernel functions.
The SGPR count includes any special extra registers such as VCC. */
@@ -98,10 +105,6 @@ gcn_init_machine_status (void)
f = ggc_cleared_alloc<machine_function> ();
- /* Set up LDS allocation for broadcasting for this function. */
- f->lds_allocated = 32;
- f->lds_allocs = hash_map<tree, int>::create_ggc (64);
-
/* And LDS temporary decls for worker reductions. */
vec_alloc (f->reduc_decls, 0);
@@ -143,6 +146,24 @@ gcn_option_override (void)
/* 1MB total. */
stack_size_opt = 1048576;
}
+
+ /* Reserve 1Kb (somewhat arbitrarily) of LDS space for reduction results and
+ worker broadcasts. */
+ if (gang_local_size_opt == -1)
+ gang_local_size_opt = 512;
+ else if (gang_local_size_opt < gangprivate_hwm)
+ gang_local_size_opt = gangprivate_hwm;
+ else if (gang_local_size_opt >= acc_lds_size - 1024)
+ {
+ /* We need some space for reductions and worker broadcasting. If the
+ user requests a large amount of gang-private LDS space, we might not
+ have enough left for the former. Increase the LDS allocation in that
+ case, although this may reduce the maximum occupancy on the
+ hardware. */
+ acc_lds_size = gang_local_size_opt + 1024;
+ if (acc_lds_size > 32768)
+ acc_lds_size = 32768;
+ }
}
/* }}} */
@@ -2959,7 +2980,7 @@ gcn_expand_prologue ()
The low-part is the address of the topmost addressable byte, which is
size-1. The high-part is an offset and should be zero. */
emit_move_insn (gen_rtx_REG (SImode, M0_REG),
- gen_int_mode (LDS_SIZE-1, SImode));
+ gen_int_mode (LDS_SIZE, SImode));
emit_insn (gen_prologue_use (gen_rtx_REG (SImode, M0_REG)));
@@ -5012,6 +5033,28 @@ gcn_fixup_accel_lto_options (tree fndecl)
}
}
+/* Implement TARGET_GOACC_SHARED_MEM_LAYOUT hook. */
+
+static void
+gcn_shared_mem_layout (unsigned HOST_WIDE_INT *lo,
+ unsigned HOST_WIDE_INT *hi,
+ int ARG_UNUSED (dims[GOMP_DIM_MAX]),
+ unsigned HOST_WIDE_INT
+ ARG_UNUSED (private_size[GOMP_DIM_MAX]),
+ unsigned HOST_WIDE_INT reduction_size[GOMP_DIM_MAX])
+{
+ *lo = gang_local_size_opt + reduction_size[GOMP_DIM_WORKER];
+ /* !!! We can maybe use dims[] to estimate the maximum number of work
+ groups/wavefronts/etc. we will launch, and therefore tune the maximum
+ amount of LDS we should use. For now, use a minimal amount to try to
+ maximise occupancy. */
+ *hi = acc_lds_size;
+ machine_function *machfun = cfun->machine;
+ machfun->reduction_base = gang_local_size_opt;
+ machfun->reduction_limit
+ = gang_local_size_opt + reduction_size[GOMP_DIM_WORKER];
+}
+
/* }}} */
/* {{{ ASM Output. */
@@ -5313,9 +5356,12 @@ gcn_section_type_flags (tree decl, const char *name, int reloc)
/* Helper function for gcn_asm_output_symbol_ref.
- FIXME: If we want to have propagation blocks allocated separately and
- statically like this, it would be better done via symbol refs and the
- assembler/linker. This is a temporary hack. */
+ FIXME: This function is used to lay out gang-private variables in LDS
+ on a per-CU basis.
+ There may be cases in which gang-local variables in different compilation
+ units could clobber each other. In that case we should be relying on the
+ linker to lay out gang-private LDS space, but that doesn't appear to be
+ possible at present. */
static void
gcn_print_lds_decl (FILE *f, tree var)
@@ -5323,7 +5369,7 @@ gcn_print_lds_decl (FILE *f, tree var)
int *offset;
machine_function *machfun = cfun->machine;
- if ((offset = machfun->lds_allocs->get (var)))
+ if ((offset = lds_allocs.get (var)))
fprintf (f, "%u", (unsigned) *offset);
else
{
@@ -5333,14 +5379,14 @@ gcn_print_lds_decl (FILE *f, tree var)
if (size > align && size > 4 && align < 8)
align = 8;
- machfun->lds_allocated = ((machfun->lds_allocated + align - 1)
- & ~(align - 1));
+ gangprivate_hwm = ((gangprivate_hwm + align - 1) & ~(align - 1));
- machfun->lds_allocs->put (var, machfun->lds_allocated);
- fprintf (f, "%u", machfun->lds_allocated);
- machfun->lds_allocated += size;
- if (machfun->lds_allocated > LDS_SIZE)
- error ("local data-share memory exhausted");
+ lds_allocs.put (var, gangprivate_hwm);
+ fprintf (f, "%u", gangprivate_hwm);
+ gangprivate_hwm += size;
+ if (gangprivate_hwm > gang_local_size_opt)
+ error ("gang-private data-share memory exhausted (increase with "
+ "-mgang-local-size=<number>)");
}
}
@@ -6323,6 +6369,8 @@ gcn_dwarf_register_span (rtx rtl)
#define TARGET_GOACC_VALIDATE_DIMS gcn_goacc_validate_dims
#undef TARGET_GOACC_WORKER_PARTITIONING
#define TARGET_GOACC_WORKER_PARTITIONING true
+#undef TARGET_GOACC_SHARED_MEM_LAYOUT
+#define TARGET_GOACC_SHARED_MEM_LAYOUT gcn_shared_mem_layout
#undef TARGET_HARD_REGNO_MODE_OK
#define TARGET_HARD_REGNO_MODE_OK gcn_hard_regno_mode_ok
#undef TARGET_HARD_REGNO_NREGS
diff --git a/gcc/config/gcn/gcn.h b/gcc/config/gcn/gcn.h
index eba4646f1bf..5ea075fcd7e 100644
--- a/gcc/config/gcn/gcn.h
+++ b/gcc/config/gcn/gcn.h
@@ -571,9 +571,8 @@ struct GTY(()) machine_function
HOST_WIDE_INT local_vars;
HOST_WIDE_INT callee_saves;
- unsigned lds_allocated;
- hash_map<tree, int> *lds_allocs;
-
+ unsigned HOST_WIDE_INT reduction_base;
+ unsigned HOST_WIDE_INT reduction_limit;
vec<tree, va_gc> *reduc_decls;
bool use_flat_addressing;
diff --git a/gcc/config/gcn/gcn.opt b/gcc/config/gcn/gcn.opt
index f23d5984a82..ce135288b21 100644
--- a/gcc/config/gcn/gcn.opt
+++ b/gcc/config/gcn/gcn.opt
@@ -73,6 +73,12 @@ mstack-size=
Target RejectNegative Joined UInteger Var(stack_size_opt) Init(-1)
-mstack-size=<number> Set the private segment size per wave-front, in bytes.
+int gang_local_size_opt = -1
+
+mgang-local-size=
+Target RejectNegative Joined UInteger Var(gang_local_size_opt) Init(-1)
+Amount of local data-share (LDS) memory to reserve for gang-local variables.
+
Wopenacc-dims
Target Var(warn_openacc_dims) Warning
Warn about invalid OpenACC dimensions.
diff --git a/gcc/config/nvptx/nvptx.c b/gcc/config/nvptx/nvptx.c
index 28f197f214a..1bf37f2ed82 100644
--- a/gcc/config/nvptx/nvptx.c
+++ b/gcc/config/nvptx/nvptx.c
@@ -6175,11 +6175,11 @@ nvptx_truly_noop_truncation (poly_uint64, poly_uint64)
/* Implement TARGET_GOACC_ADJUST_PRIVATE_DECL. Set "oacc gangprivate"
attribute for gang-private variable declarations. */
-void
+static tree
nvptx_goacc_adjust_private_decl (tree decl, int level)
{
if (level != GOMP_DIM_GANG)
- return;
+ return decl;
if (!lookup_attribute ("oacc gangprivate", DECL_ATTRIBUTES (decl)))
{
@@ -6192,6 +6192,8 @@ nvptx_goacc_adjust_private_decl (tree decl, int level)
tree id = get_identifier ("oacc gangprivate");
DECL_ATTRIBUTES (decl) = tree_cons (id, NULL, DECL_ATTRIBUTES (decl));
}
+
+ return decl;
}
/* Implement TARGET_GOACC_EXPAND_ACCEL_VAR. Place "oacc gangprivate"
diff --git a/gcc/doc/tm.texi b/gcc/doc/tm.texi
index 54acc5169b6..8259b6dbb38 100644
--- a/gcc/doc/tm.texi
+++ b/gcc/doc/tm.texi
@@ -6235,7 +6235,7 @@ memories. A return value of NULL indicates that the target does not
handle this VAR_DECL, and normal RTL expanding is resumed.
@end deftypefn
-@deftypefn {Target Hook} void TARGET_GOACC_ADJUST_PRIVATE_DECL (tree @var{var}, @var{int})
+@deftypefn {Target Hook} tree TARGET_GOACC_ADJUST_PRIVATE_DECL (tree @var{var}, @var{int})
Tweak variable declaration for a private variable at the specified
parallelism level.
@end deftypefn
@@ -6244,12 +6244,19 @@ parallelism level.
Use gimple transformation for worker neutering/broadcasting.
@end deftypevr
-@deftypefn {Target Hook} tree TARGET_GOACC_CREATE_PROPAGATION_RECORD (tree @var{rec}, bool @var{sender}, const char *@var{name})
+@deftypefn {Target Hook} tree TARGET_GOACC_CREATE_PROPAGATION_RECORD (tree @var{rec}, bool @var{sender}, const char *@var{name}, unsigned HOST_WIDE_INT @var{offset})
Create a record used to propagate local-variable state from an active
worker to other workers. A possible implementation might adjust the type
of REC to place the new variable in shared GPU memory.
@end deftypefn
+@deftypefn {Target Hook} void TARGET_GOACC_SHARED_MEM_LAYOUT (unsigned HOST_WIDE_INT *@var{}, unsigned HOST_WIDE_INT *@var{}, @var{int[]}, unsigned @var{HOST_WIDE_INT[]}, unsigned @var{HOST_WIDE_INT[]})
+Lay out a fixed shared-memory region on the target. The LO and HI
+arguments should be set to a range of addresses that can be used for worker
+broadcasting. The dimensions, reduction size and gang-private size
+arguments are for the current offload region.
+@end deftypefn
+
@node Anchored Addresses
@section Anchored Addresses
@cindex anchored addresses
diff --git a/gcc/doc/tm.texi.in b/gcc/doc/tm.texi.in
index 6c97edecb2e..193b3478b5c 100644
--- a/gcc/doc/tm.texi.in
+++ b/gcc/doc/tm.texi.in
@@ -4227,6 +4227,8 @@ address; but often a machine-dependent strategy can generate better code.
@hook TARGET_GOACC_CREATE_PROPAGATION_RECORD
+@hook TARGET_GOACC_SHARED_MEM_LAYOUT
+
@node Anchored Addresses
@section Anchored Addresses
@cindex anchored addresses
diff --git a/gcc/omp-offload.c b/gcc/omp-offload.c
index 22bb3ac1c40..dec9bdfd5b1 100644
--- a/gcc/omp-offload.c
+++ b/gcc/omp-offload.c
@@ -1847,7 +1847,7 @@ maybe_discard_oacc_function (tree decl)
struct addr_expr_rewrite_info
{
gimple *stmt;
- hash_set<tree> *adjusted_vars;
+ hash_map<tree, tree> *adjusted_vars;
bool avoid_pointer_conversion;
bool modified;
};
@@ -1861,19 +1861,20 @@ rewrite_addr_expr (tree *tp, int *walk_subtrees, void *data)
if (TREE_CODE (*tp) == ADDR_EXPR)
{
tree arg = TREE_OPERAND (*tp, 0);
+ tree *new_arg = info->adjusted_vars->get (arg);
- if (info->adjusted_vars->contains (arg))
+ if (new_arg)
{
if (info->avoid_pointer_conversion)
{
- *tp = build_fold_addr_expr (arg);
+ *tp = build_fold_addr_expr (*new_arg);
info->modified = true;
*walk_subtrees = 0;
}
else
{
gimple_stmt_iterator gsi = gsi_for_stmt (info->stmt);
- tree repl = build_fold_addr_expr (arg);
+ tree repl = build_fold_addr_expr (*new_arg);
gimple *stmt1
= gimple_build_assign (make_ssa_name (TREE_TYPE (repl)), repl);
tree conv = convert_to_pointer (TREE_TYPE (*tp),
@@ -1888,6 +1889,15 @@ rewrite_addr_expr (tree *tp, int *walk_subtrees, void *data)
}
}
}
+ else if (TREE_CODE (*tp) == VAR_DECL)
+ {
+ tree *new_decl = info->adjusted_vars->get (*tp);
+ if (new_decl)
+ {
+ *tp = *new_decl;
+ info->modified = true;
+ }
+ }
return NULL_TREE;
}
@@ -1921,7 +1931,8 @@ is_sync_builtin_call (gcall *call)
tree
default_goacc_create_propagation_record (tree record_type, bool sender,
- const char *name)
+ const char *name,
+ unsigned HOST_WIDE_INT ARG_UNUSED (offset))
{
tree type = record_type;
@@ -2096,8 +2107,106 @@ execute_oacc_loop_designation ()
int
execute_oacc_gimple_workers (void)
{
- oacc_do_neutering ();
- calculate_dominance_info (CDI_DOMINATORS);
+ unsigned HOST_WIDE_INT reduction_size[GOMP_DIM_MAX];
+ unsigned HOST_WIDE_INT private_size[GOMP_DIM_MAX];
+
+ for (unsigned i = 0; i < GOMP_DIM_MAX; i++)
+ {
+ reduction_size[i] = 0;
+ private_size[i] = 0;
+ }
+
+ /* Calculate shared memory size required for reduction variables and
+ gang-private memory for this offloaded function. */
+ basic_block bb;
+ FOR_ALL_BB_FN (bb, cfun)
+ {
+ for (gimple_stmt_iterator gsi = gsi_start_bb (bb);
+ !gsi_end_p (gsi);
+ gsi_next (&gsi))
+ {
+ gimple *stmt = gsi_stmt (gsi);
+ if (!is_gimple_call (stmt))
+ continue;
+ gcall *call = as_a <gcall *> (stmt);
+ if (!gimple_call_internal_p (call))
+ continue;
+ enum internal_fn ifn_code = gimple_call_internal_fn (call);
+ switch (ifn_code)
+ {
+ default: break;
+ case IFN_GOACC_REDUCTION:
+ if (integer_minus_onep (gimple_call_arg (call, 3)))
+ continue;
+ else
+ {
+ unsigned code = TREE_INT_CST_LOW (gimple_call_arg (call, 0));
+ /* Only count reduction variables once: the choice to pick
+ the setup call is fairly arbitrary. */
+ if (code == IFN_GOACC_REDUCTION_SETUP)
+ {
+ int level = TREE_INT_CST_LOW (gimple_call_arg (call, 3));
+ tree var = gimple_call_arg (call, 2);
+ tree offset = gimple_call_arg (call, 5);
+ tree var_type = TREE_TYPE (var);
+ unsigned HOST_WIDE_INT limit
+ = tree_to_uhwi (offset)
+ + tree_to_uhwi (TYPE_SIZE_UNIT (var_type));
+ reduction_size[level]
+ = MAX (reduction_size[level], limit);
+ }
+ }
+ break;
+ case IFN_UNIQUE:
+ {
+ enum ifn_unique_kind kind
+ = ((enum ifn_unique_kind)
+ TREE_INT_CST_LOW (gimple_call_arg (call, 0)));
+
+ if (kind == IFN_UNIQUE_OACC_PRIVATE)
+ {
+ HOST_WIDE_INT level
+ = TREE_INT_CST_LOW (gimple_call_arg (call, 2));
+ if (level == -1)
+ break;
+ for (unsigned i = 3;
+ i < gimple_call_num_args (call);
+ i++)
+ {
+ tree arg = gimple_call_arg (call, i);
+ gcc_assert (TREE_CODE (arg) == ADDR_EXPR);
+ tree decl = TREE_OPERAND (arg, 0);
+ unsigned HOST_WIDE_INT align = DECL_ALIGN_UNIT (decl);
+ private_size[level] = ((private_size[level] + align - 1)
+ & ~(align - 1));
+ unsigned HOST_WIDE_INT decl_size
+ = tree_to_uhwi (TYPE_SIZE_UNIT (TREE_TYPE (decl)));
+ private_size[level] += decl_size;
+ }
+ }
+ }
+ break;
+ }
+ }
+ }
+
+ int dims[GOMP_DIM_MAX];
+
+ for (unsigned i = 0; i < GOMP_DIM_MAX; i++)
+ dims[i] = oacc_get_fn_dim_size (current_function_decl, i);
+
+ /* Find bounds of shared-memory buffer space we can use. */
+ unsigned HOST_WIDE_INT bounds_lo = 0, bounds_hi = 0;
+ if (targetm.goacc.shared_mem_layout)
+ targetm.goacc.shared_mem_layout (&bounds_lo, &bounds_hi, dims,
+ private_size, reduction_size);
+
+ /* Perform worker partitioning unless we know the number of workers is 1. */
+ if (dims[GOMP_DIM_WORKER] != 1)
+ {
+ oacc_do_neutering (bounds_lo, bounds_hi);
+ calculate_dominance_info (CDI_DOMINATORS);
+ }
return 0;
}
@@ -2114,7 +2223,7 @@ execute_oacc_device_lower ()
for (unsigned i = 0; i < GOMP_DIM_MAX; i++)
dims[i] = oacc_get_fn_dim_size (current_function_decl, i);
- hash_set<tree> adjusted_vars;
+ hash_map<tree, tree> adjusted_vars;
/* Now lower internal loop functions to target-specific code
sequences. */
@@ -2221,9 +2330,11 @@ execute_oacc_device_lower ()
if (targetm.goacc.adjust_private_decl)
{
tree oldtype = TREE_TYPE (decl);
- targetm.goacc.adjust_private_decl (decl, level);
- if (TREE_TYPE (decl) != oldtype)
- adjusted_vars.add (decl);
+ tree newdecl
+ = targetm.goacc.adjust_private_decl (decl, level);
+ if (TREE_TYPE (newdecl) != oldtype
+ || newdecl != decl)
+ adjusted_vars.put (decl, newdecl);
}
}
remove = true;
@@ -2398,22 +2509,11 @@ public:
{}
/* opt_pass methods: */
- virtual bool gate (function *)
+ virtual bool gate (function *fun)
{
- if (!flag_openacc || !targetm.goacc.worker_partitioning)
- return false;
-
- tree attr = oacc_get_fn_attrib (current_function_decl);
-
- if (!attr)
- /* Not an offloaded function. */
- return false;
-
- int worker_dim
- = oacc_get_fn_dim_size (current_function_decl, GOMP_DIM_WORKER);
-
- /* No worker partitioning if we know the number of workers is 1. */
- return worker_dim != 1;
+ return flag_openacc
+ && targetm.goacc.worker_partitioning
+ && oacc_get_fn_attrib (fun->decl);
};
virtual unsigned int execute (function *)
diff --git a/gcc/omp-sese.c b/gcc/omp-sese.c
index c5d87ab729a..c279cf1708a 100644
--- a/gcc/omp-sese.c
+++ b/gcc/omp-sese.c
@@ -53,6 +53,8 @@
#include "omp-offload.h"
#include "attribs.h"
#include "omp-sese.h"
+#include "targhooks.h"
+#include "diagnostic-core.h"
/* Loop structure of the function. The entire function is described as
a NULL loop. */
@@ -966,6 +968,8 @@ static tree
build_sender_ref (tree record_type, tree var, tree sender_decl)
{
field_map_t *fields = *field_map->get (record_type);
+ if (POINTER_TYPE_P (TREE_TYPE (sender_decl)))
+ sender_decl = build_simple_mem_ref (sender_decl);
tree field = *fields->get (var);
return oacc_build_component_ref (sender_decl, field);
}
@@ -1003,7 +1007,8 @@ static void
worker_single_copy (basic_block from, basic_block to,
hash_set<tree> *def_escapes_block,
hash_set<tree> *worker_partitioned_uses,
- tree record_type)
+ tree record_type, unsigned HOST_WIDE_INT placement,
+ bool isolate_broadcasts)
{
/* If we only have virtual defs, we'll have no record type, but we still want
to emit single_copy_start and (particularly) single_copy_end to act as
@@ -1014,10 +1019,10 @@ worker_single_copy (basic_block from, basic_block to,
tree sender_decl
= targetm.goacc.create_propagation_record (record_type, true,
- ".oacc_worker_o");
+ ".oacc_worker_o", placement);
tree receiver_decl
= targetm.goacc.create_propagation_record (record_type, false,
- ".oacc_worker_i");
+ ".oacc_worker_i", placement);
gimple_stmt_iterator gsi = gsi_last_bb (to);
if (EDGE_COUNT (to->succs) > 1)
@@ -1031,12 +1036,23 @@ worker_single_copy (basic_block from, basic_block to,
tree lhs = create_tmp_var (TREE_TYPE (TREE_TYPE (decl)));
- gimple *call = gimple_build_call (decl, 1,
- build_fold_addr_expr (sender_decl));
+ gimple *call
+ = gimple_build_call (decl, 1,
+ POINTER_TYPE_P (TREE_TYPE (sender_decl))
+ ? sender_decl : build_fold_addr_expr (sender_decl));
gimple_call_set_lhs (call, lhs);
gsi_insert_before (&start, call, GSI_NEW_STMT);
update_stmt (call);
+ /* The shared-memory range for this block overflowed. Add a barrier before
+ the GOACC_single_copy_start call. */
+ if (isolate_broadcasts)
+ {
+ decl = builtin_decl_explicit (BUILT_IN_GOACC_BARRIER);
+ gimple *acc_bar = gimple_build_call (decl, 0);
+ gsi_insert_before (&start, acc_bar, GSI_SAME_STMT);
+ }
+
tree conv_tmp = make_ssa_name (TREE_TYPE (receiver_decl));
gimple *conv = gimple_build_assign (conv_tmp,
@@ -1205,13 +1221,26 @@ worker_single_copy (basic_block from, basic_block to,
}
}
+ /* The shared-memory range for this block overflowed. Add a barrier at the
+ end. */
+ if (isolate_broadcasts)
+ {
+ gsi = gsi_start_bb (exit_block);
+ decl = builtin_decl_explicit (BUILT_IN_GOACC_BARRIER);
+ gimple *acc_bar = gimple_build_call (decl, 0);
+ gsi_insert_before (&gsi, acc_bar, GSI_SAME_STMT);
+ }
+
/* It's possible for the ET->DEST block (the work done by the active thread)
to finish with a control-flow insn, e.g. a UNIQUE function call. Split
the block and add SENDER_SEQ in the latter part to avoid having control
flow in the middle of a BB. */
decl = builtin_decl_explicit (BUILT_IN_GOACC_SINGLE_COPY_END);
- call = gimple_build_call (decl, 1, build_fold_addr_expr (sender_decl));
+ call = gimple_build_call (decl, 1,
+ POINTER_TYPE_P (TREE_TYPE (sender_decl))
+ ? sender_decl
+ : build_fold_addr_expr (sender_decl));
gimple_seq_add_stmt (&sender_seq, call);
gsi = gsi_last_bb (body);
@@ -1221,11 +1250,15 @@ worker_single_copy (basic_block from, basic_block to,
gsi_insert_seq_after (&gsi, sender_seq, GSI_CONTINUE_LINKING);
}
+typedef hash_map<basic_block, std::pair<unsigned HOST_WIDE_INT, bool> >
+ blk_offset_map_t;
+
static void
neuter_worker_single (parallel_g *par, unsigned outer_mask,
bitmap worker_single, bitmap vector_single,
vec<propagation_set *> *prop_set,
- hash_set<tree> *partitioned_var_uses)
+ hash_set<tree> *partitioned_var_uses,
+ blk_offset_map_t *blk_offset_map)
{
unsigned mask = outer_mask | par->mask;
@@ -1314,8 +1347,18 @@ neuter_worker_single (parallel_g *par, unsigned outer_mask,
tree record_type = (tree) block->aux;
if (has_defs)
- worker_single_copy (block, block, &def_escapes_block,
- &worker_partitioned_uses, record_type);
+ {
+ std::pair<unsigned HOST_WIDE_INT, bool> *off_rngalloc
+ = blk_offset_map->get (block);
+ gcc_assert (!record_type || off_rngalloc);
+ unsigned HOST_WIDE_INT offset
+ = off_rngalloc ? off_rngalloc->first : 0;
+ bool range_allocated
+ = off_rngalloc ? off_rngalloc->second : true;
+ worker_single_copy (block, block, &def_escapes_block,
+ &worker_partitioned_uses, record_type,
+ offset, !range_allocated);
+ }
else
worker_single_simple (block, block, &def_escapes_block);
}
@@ -1351,15 +1394,158 @@ neuter_worker_single (parallel_g *par, unsigned outer_mask,
if (par->inner)
neuter_worker_single (par->inner, mask, worker_single, vector_single,
- prop_set, partitioned_var_uses);
+ prop_set, partitioned_var_uses, blk_offset_map);
if (par->next)
neuter_worker_single (par->next, outer_mask, worker_single, vector_single,
- prop_set, partitioned_var_uses);
+ prop_set, partitioned_var_uses, blk_offset_map);
+}
+
+
+static void
+dfs_broadcast_reachable_1 (basic_block bb, sbitmap reachable)
+{
+ if (bb->flags & BB_VISITED)
+ return;
+
+ bb->flags |= BB_VISITED;
+
+ if (bb->succs)
+ {
+ edge e;
+ edge_iterator ei;
+ FOR_EACH_EDGE (e, ei, bb->succs)
+ {
+ basic_block dest = e->dest;
+ if (dest->aux)
+ bitmap_set_bit (reachable, dest->index);
+ else
+ dfs_broadcast_reachable_1 (dest, reachable);
+ }
+ }
+}
+
+typedef std::pair<int, tree> idx_decl_pair_t;
+
+typedef auto_vec<splay_tree> used_range_vec_t;
+
+static int
+sort_size_descending (const void *a, const void *b)
+{
+ const idx_decl_pair_t *pa = (const idx_decl_pair_t *) a;
+ const idx_decl_pair_t *pb = (const idx_decl_pair_t *) b;
+ unsigned HOST_WIDE_INT asize = tree_to_uhwi (TYPE_SIZE_UNIT (pa->second));
+ unsigned HOST_WIDE_INT bsize = tree_to_uhwi (TYPE_SIZE_UNIT (pb->second));
+ return bsize - asize;
+}
+
+class addr_range
+{
+public:
+ addr_range (unsigned HOST_WIDE_INT addr_lo, unsigned HOST_WIDE_INT addr_hi)
+ : lo (addr_lo), hi (addr_hi)
+ { }
+ addr_range (const addr_range &ar) : lo (ar.lo), hi (ar.hi)
+ { }
+ addr_range () : lo (0), hi (0)
+ { }
+
+ bool invalid () { return lo == 0 && hi == 0; }
+
+ unsigned HOST_WIDE_INT lo;
+ unsigned HOST_WIDE_INT hi;
+};
+
+static int
+splay_tree_compare_addr_range (splay_tree_key a, splay_tree_key b)
+{
+ addr_range *ar = (addr_range *) a;
+ addr_range *br = (addr_range *) b;
+ if (ar->lo == br->lo && ar->hi == br->hi)
+ return 0;
+ if (ar->hi <= br->lo)
+ return -1;
+ else if (ar->lo >= br->hi)
+ return 1;
+ return 0;
+}
+
+static void
+splay_tree_free_key (splay_tree_key k)
+{
+ addr_range *ar = (addr_range *) k;
+ delete ar;
}
+static addr_range
+first_fit_range (splay_tree s, unsigned HOST_WIDE_INT size,
+ unsigned HOST_WIDE_INT align, addr_range *bounds)
+{
+ splay_tree_node min = splay_tree_min (s);
+ if (min)
+ {
+ splay_tree_node next;
+ while ((next = splay_tree_successor (s, min->key)))
+ {
+ unsigned HOST_WIDE_INT lo = ((addr_range *) min->key)->hi;
+ unsigned HOST_WIDE_INT hi = ((addr_range *) next->key)->lo;
+ unsigned HOST_WIDE_INT base = (lo + align - 1) & ~(align - 1);
+ if (base + size <= hi)
+ return addr_range (base, base + size);
+ min = next;
+ }
+
+ unsigned HOST_WIDE_INT base = ((addr_range *)min->key)->hi;
+ base = (base + align - 1) & ~(align - 1);
+ if (base + size <= bounds->hi)
+ return addr_range (base, base + size);
+ else
+ return addr_range ();
+ }
+ else
+ {
+ unsigned HOST_WIDE_INT lo = bounds->lo;
+ lo = (lo + align - 1) & ~(align - 1);
+ if (lo + size <= bounds->hi)
+ return addr_range (lo, lo + size);
+ else
+ return addr_range ();
+ }
+}
+
+static int
+merge_ranges_1 (splay_tree_node n, void *ptr)
+{
+ splay_tree accum = (splay_tree) ptr;
+ addr_range ar = *(addr_range *) n->key;
+
+ splay_tree_node old = splay_tree_lookup (accum, n->key);
+
+ /* We might have an overlap. Create a new range covering the
+ overlapping parts. */
+ if (old)
+ {
+ addr_range *old_ar = (addr_range *) old->key;
+ ar.lo = MIN (old_ar->lo, ar.lo);
+ ar.hi = MAX (old_ar->hi, ar.hi);
+ splay_tree_remove (accum, old->key);
+ }
+
+ addr_range *new_ar = new addr_range (ar);
+
+ splay_tree_insert (accum, (splay_tree_key) new_ar, n->value);
+
+ return 0;
+}
+
+static void
+merge_ranges (splay_tree accum, splay_tree sp)
+{
+ splay_tree_foreach (sp, merge_ranges_1, (void *) accum);
+}
void
-oacc_do_neutering (void)
+oacc_do_neutering (unsigned HOST_WIDE_INT bounds_lo,
+ unsigned HOST_WIDE_INT bounds_hi)
{
bb_stmt_map_t bb_stmt_map;
auto_bitmap worker_single, vector_single;
@@ -1448,8 +1634,122 @@ oacc_do_neutering (void)
}
}
+ sbitmap *reachable
+ = sbitmap_vector_alloc (last_basic_block_for_fn (cfun),
+ last_basic_block_for_fn (cfun));
+
+ bitmap_vector_clear (reachable, last_basic_block_for_fn (cfun));
+
+ auto_vec<std::pair<int, tree> > priority;
+
+ FOR_ALL_BB_FN (bb, cfun)
+ {
+ if (bb->aux)
+ {
+ tree record_type = (tree) bb->aux;
+
+ basic_block bb2;
+ FOR_ALL_BB_FN (bb2, cfun)
+ bb2->flags &= ~BB_VISITED;
+
+ priority.safe_push (std::make_pair (bb->index, record_type));
+ dfs_broadcast_reachable_1 (bb, reachable[bb->index]);
+ }
+ }
+
+ sbitmap *inverted
+ = sbitmap_vector_alloc (last_basic_block_for_fn (cfun),
+ last_basic_block_for_fn (cfun));
+
+ bitmap_vector_clear (inverted, last_basic_block_for_fn (cfun));
+
+ for (int i = 0; i < last_basic_block_for_fn (cfun); i++)
+ {
+ sbitmap_iterator bi;
+ unsigned int j;
+ EXECUTE_IF_SET_IN_BITMAP (reachable[i], 0, j, bi)
+ bitmap_set_bit (inverted[j], i);
+ }
+
+ for (int i = 0; i < last_basic_block_for_fn (cfun); i++)
+ bitmap_ior (reachable[i], reachable[i], inverted[i]);
+
+ sbitmap_vector_free (inverted);
+
+ used_range_vec_t used_ranges;
+
+ used_ranges.safe_grow_cleared (last_basic_block_for_fn (cfun));
+
+ blk_offset_map_t blk_offset_map;
+
+ addr_range worker_shm_bounds (bounds_lo, bounds_hi);
+
+ priority.qsort (sort_size_descending);
+ for (unsigned int i = 0; i < priority.length (); i++)
+ {
+ idx_decl_pair_t p = priority[i];
+ int blkno = p.first;
+ tree record_type = p.second;
+ HOST_WIDE_INT size = tree_to_uhwi (TYPE_SIZE_UNIT (record_type));
+ HOST_WIDE_INT align = TYPE_ALIGN_UNIT (record_type);
+
+ splay_tree conflicts = splay_tree_new (splay_tree_compare_addr_range,
+ splay_tree_free_key, NULL);
+
+ if (!used_ranges[blkno])
+ used_ranges[blkno] = splay_tree_new (splay_tree_compare_addr_range,
+ splay_tree_free_key, NULL);
+ else
+ merge_ranges (conflicts, used_ranges[blkno]);
+
+ sbitmap_iterator bi;
+ unsigned int j;
+ EXECUTE_IF_SET_IN_BITMAP (reachable[blkno], 0, j, bi)
+ if (used_ranges[j])
+ merge_ranges (conflicts, used_ranges[j]);
+
+ addr_range ar
+ = first_fit_range (conflicts, size, align, &worker_shm_bounds);
+
+ splay_tree_delete (conflicts);
+
+ if (ar.invalid ())
+ {
+ unsigned HOST_WIDE_INT base;
+ base = bounds_lo + random () % 512;
+ base = (base + align - 1) & ~(align - 1);
+ if (base + size > bounds_hi)
+ error_at (UNKNOWN_LOCATION, "shared-memory region overflow");
+ std::pair<unsigned HOST_WIDE_INT, bool> base_inrng
+ = std::make_pair (base, false);
+ blk_offset_map.put (BASIC_BLOCK_FOR_FN (cfun, blkno), base_inrng);
+ }
+ else
+ {
+ splay_tree_node old = splay_tree_lookup (used_ranges[blkno],
+ (splay_tree_key) &ar);
+ if (old)
+ {
+ fprintf (stderr, "trying to map [%d..%d] but [%d..%d] is "
+ "already mapped in block %d\n", (int) ar.lo,
+ (int) ar.hi, (int) ((addr_range *) old->key)->lo,
+ (int) ((addr_range *) old->key)->hi, blkno);
+ abort ();
+ }
+
+ addr_range *arp = new addr_range (ar);
+ splay_tree_insert (used_ranges[blkno], (splay_tree_key) arp,
+ (splay_tree_value) blkno);
+ std::pair<unsigned HOST_WIDE_INT, bool> base_inrng
+ = std::make_pair (ar.lo, true);
+ blk_offset_map.put (BASIC_BLOCK_FOR_FN (cfun, blkno), base_inrng);
+ }
+ }
+
+ sbitmap_vector_free (reachable);
+
neuter_worker_single (par, mask, worker_single, vector_single, &prop_set,
- &partitioned_var_uses);
+ &partitioned_var_uses, &blk_offset_map);
prop_set.release ();
diff --git a/gcc/omp-sese.h b/gcc/omp-sese.h
index 0b82a2417eb..9aa25e8616e 100644
--- a/gcc/omp-sese.h
+++ b/gcc/omp-sese.h
@@ -27,6 +27,6 @@ typedef auto_vec<bb_pair_t> bb_pair_vec_t;
extern void omp_find_sese (auto_vec<basic_block> &blocks,
bb_pair_vec_t ®ions);
-extern void oacc_do_neutering (void);
+extern void oacc_do_neutering (unsigned HOST_WIDE_INT, unsigned HOST_WIDE_INT);
#endif
diff --git a/gcc/target.def b/gcc/target.def
index a3f67fa8223..3b3719dd3b7 100644
--- a/gcc/target.def
+++ b/gcc/target.def
@@ -1740,7 +1740,7 @@ DEFHOOK
(adjust_private_decl,
"Tweak variable declaration for a private variable at the specified\n\
parallelism level.",
-void, (tree var, int),
+tree, (tree var, int),
NULL)
DEFHOOK
@@ -1748,7 +1748,7 @@ DEFHOOK
"Create a record used to propagate local-variable state from an active\n\
worker to other workers. A possible implementation might adjust the type\n\
of REC to place the new variable in shared GPU memory.",
-tree, (tree rec, bool sender, const char *name),
+tree, (tree rec, bool sender, const char *name, unsigned HOST_WIDE_INT offset),
default_goacc_create_propagation_record)
DEFHOOKPOD
@@ -1756,6 +1756,16 @@ DEFHOOKPOD
"Use gimple transformation for worker neutering/broadcasting.",
bool, false)
+DEFHOOK
+(shared_mem_layout,
+"Lay out a fixed shared-memory region on the target. The LO and HI\n\
+arguments should be set to a range of addresses that can be used for worker\n\
+broadcasting. The dimensions, reduction size and gang-private size\n\
+arguments are for the current offload region.",
+void, (unsigned HOST_WIDE_INT *, unsigned HOST_WIDE_INT *, int[],
+ unsigned HOST_WIDE_INT[], unsigned HOST_WIDE_INT[]),
+NULL)
+
HOOK_VECTOR_END (goacc)
/* Functions relating to vectorization. */
diff --git a/gcc/targhooks.h b/gcc/targhooks.h
index 281fa39d39a..0c1e9a2c3db 100644
--- a/gcc/targhooks.h
+++ b/gcc/targhooks.h
@@ -130,7 +130,8 @@ extern bool default_goacc_validate_dims (tree, int [], int, unsigned);
extern int default_goacc_dim_limit (int);
extern bool default_goacc_fork_join (gcall *, const int [], bool);
extern void default_goacc_reduction (gcall *);
-extern tree default_goacc_create_propagation_record (tree, bool, const char *);
+extern tree default_goacc_create_propagation_record (tree, bool, const char *,
+ unsigned HOST_WIDE_INT);
/* These are here, and not in hooks.[ch], because not all users of
hooks.h include tm.h, and thus we don't have CUMULATIVE_ARGS. */
diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp
index 48be697a23e..30a8f71906d 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,7 @@
+2020-07-15 Julian Brown <julian@codesourcery.com>
+
+ * testsuite/libgomp.oacc-c-c++-common/broadcast-many.c: New test.
+
2020-07-15 Julian Brown <julian@codesourcery.com>
* testsuite/libgomp.oacc-fortran/collapse-1.f90: Fix race condition.
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-many.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-many.c
new file mode 100644
index 00000000000..70caea81dd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-many.c
@@ -0,0 +1,79 @@
+/* { dg-additional-options "-foffload=-mgang-local-size=64" { target openacc_radeon_accel_selected } } */
+
+#include <assert.h>
+#include <stdio.h>
+
+#define LOCAL(n) double n = input;
+#define LOCALS(n) LOCAL(n##1) LOCAL(n##2) LOCAL(n##3) LOCAL(n##4) \
+ LOCAL(n##5) LOCAL(n##6) LOCAL(n##7) LOCAL(n##8)
+#define LOCALS2(n) LOCALS(n##a) LOCALS(n##b) LOCALS(n##c) LOCALS(n##d) \
+ LOCALS(n##e) LOCALS(n##f) LOCALS(n##g) LOCALS(n##h)
+
+#define USE(n) n
+#define USES(n,OP) USE(n##1) OP USE(n##2) OP USE(n##3) OP USE (n##4) OP \
+ USE(n##5) OP USE(n##6) OP USE(n##7) OP USE (n##8)
+#define USES2(n,OP) USES(n##a,OP) OP USES(n##b,OP) OP USES(n##c,OP) OP \
+ USES(n##d,OP) OP USES(n##e,OP) OP USES(n##f,OP) OP \
+ USES(n##g,OP) OP USES(n##h,OP)
+
+int main (void)
+{
+ int ret;
+ int input = 1;
+
+ #pragma acc parallel num_gangs(1) num_workers(32) copyout(ret)
+ {
+ int w = 0;
+ LOCALS2(h);
+
+ #pragma acc loop worker reduction(+:w)
+ for (int i = 0; i < 32; i++)
+ {
+ int u = USES2(h,+);
+ w += u;
+ }
+
+ printf ("w=%d\n", w);
+ /* { dg-output "w=2048(\n|\r\n|\r)" } */
+
+ LOCALS2(i);
+
+ #pragma acc loop worker reduction(+:w)
+ for (int i = 0; i < 32; i++)
+ {
+ int u = USES2(i,+);
+ w += u;
+ }
+
+ printf ("w=%d\n", w);
+ /* { dg-output "w=4096(\n|\r\n|\r)" } */
+
+ LOCALS2(j);
+ LOCALS2(k);
+
+ #pragma acc loop worker reduction(+:w)
+ for (int i = 0; i < 32; i++)
+ {
+ int u = USES2(j,+);
+ w += u;
+ }
+
+ printf ("w=%d\n", w);
+ /* { dg-output "w=6144(\n|\r\n|\r)" } */
+
+ #pragma acc loop worker reduction(+:w)
+ for (int i = 0; i < 32; i++)
+ {
+ int u = USES2(k,+);
+ w += u;
+ }
+
+ ret = (w == 64 * 32 * 4);
+ printf ("w=%d\n", w);
+ /* { dg-output "w=8192(\n|\r\n|\r)" } */
+ }
+
+ assert (ret);
+
+ return 0;
+}
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:15 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:15 [gcc/devel/omp/gcc-11] openacc: Shared memory layout optimisation Kwok Yeung
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).