From patchwork Mon Jun 29 20:16:52 2020 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1319164 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces@gcc.gnu.org; receiver=) Authentication-Results: ozlabs.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by ozlabs.org (Postfix) with ESMTPS id 49wf1F42Drz9sR4 for ; Tue, 30 Jun 2020 06:17:37 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 9C0633840C25; Mon, 29 Jun 2020 20:17:28 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 27F223870868 for ; Mon, 29 Jun 2020 20:17:23 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.3.2 sourceware.org 27F223870868 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=Julian_Brown@mentor.com IronPort-SDR: JVEFYVVKfcBks7xGfCwQKC+9A4sbuS+VD84S7qNpLYWzZMiAap3d2K3blziocr0VMHmSX9ECKP y0VZpapHSNtUvGp5inJ7g6Ek8amjjxeKmIxNlNCWuxtqRNUGt581jkmtGU4wt8HMCiDWhhddG7 QiQff2trWRc59JrNJT4R35VddxOQUU5O0SnaUfTReC2ni9WgkTrJ4+ITJwGzXiJNxPvo5/vMJL cOvuqEXho3ez7PcpH6wRvpDmRdh+nSUAFe0T7apLo/VJ/+HyeL4q7xjwAUVaba14eikaK5nE0r FqI= X-IronPort-AV: E=Sophos;i="5.75,295,1589270400"; d="scan'208";a="50429885" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa2.mentor.iphmx.com with ESMTP; 29 Jun 2020 12:17:22 -0800 IronPort-SDR: pCkUozTSw3hGlTmJ/CPUooF9r7yiIDe+jnn0S8OlUJTG1ZnLPCZDjuHwEwEHHXQRNpyp8J7ju/ i11Un4yrq4DkBXwEr3+0fHkcB91/MfNdNMrFVlEeob3j79lqvRWMXr3i9Rmys/mrW9A3rtABiU ItCTNu3N/mi+tobnSvW9Hy5M/rYdvytQVmuDmJn6sNDTS3SPUaWO8TMRS8jyBmVlbHobTZwpsM sdmXafIo2HA92nNwpTC87pMADCiTzZbnyGWroWVMcWIlXXhqVgSBA7XUwhGRFB9QkznZKKwAQG R1c= From: Julian Brown To: Subject: [PATCH] [og10] OpenACC: Shared memory layout optimisation Date: Mon, 29 Jun 2020 13:16:52 -0700 Message-ID: <20200629201653.147326-4-julian@codesourcery.com> X-Mailer: git-send-email 2.23.0 In-Reply-To: <20200629201653.147326-1-julian@codesourcery.com> References: <20200629201653.147326-1-julian@codesourcery.com> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-08.mgc.mentorg.com (139.181.222.8) To SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) X-Spam-Status: No, score=-12.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.2 X-Spam-Checker-Version: SpamAssassin 3.4.2 (2018-09-13) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Jakub Jelinek , Andrew Stubbs Errors-To: gcc-patches-bounces@gcc.gnu.org Sender: "Gcc-patches" 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= (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. OK for og10 branch? Julian ChangeLog 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. * 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. --- 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/doc/tm.texi | 11 +- gcc/doc/tm.texi.in | 2 + gcc/omp-offload.c | 148 ++++++-- gcc/omp-sese.c | 323 +++++++++++++++++- gcc/omp-sese.h | 2 +- gcc/target.def | 14 +- gcc/targhooks.h | 3 +- .../broadcast-many.c | 79 +++++ 13 files changed, 638 insertions(+), 90 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/broadcast-many.c diff --git a/gcc/config/gcn/gcn-protos.h b/gcc/config/gcn/gcn-protos.h index b728c0463df1..5bd9f44572f8 100644 --- a/gcc/config/gcn/gcn-protos.h +++ b/gcc/config/gcn/gcn-protos.h @@ -38,8 +38,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 e45dc047be34..1dc9b21823b0 100644 --- a/gcc/config/gcn/gcn-tree.c +++ b/gcc/config/gcn/gcn-tree.c @@ -448,14 +448,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); (*machfun->reduc_decls)[offset] = decl; @@ -674,34 +672,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, @@ -720,6 +720,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 f9848a26dc7c..bf996b461547 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 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 (); - /* Set up LDS allocation for broadcasting for this function. */ - f->lds_allocated = 32; - f->lds_allocs = hash_map::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; + } } /* }}} */ @@ -2888,7 +2909,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))); @@ -4920,6 +4941,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. */ @@ -5220,9 +5263,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) @@ -5230,7 +5276,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 { @@ -5240,14 +5286,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=)"); } } @@ -6170,6 +6216,8 @@ print_operand (FILE *file, rtx x, int code) #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 f63e7df3021a..91b3d88c062a 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 *lds_allocs; - + unsigned HOST_WIDE_INT reduction_base; + unsigned HOST_WIDE_INT reduction_limit; vec *reduc_decls; bool use_flat_addressing; diff --git a/gcc/config/gcn/gcn.opt b/gcc/config/gcn/gcn.opt index 8e4a7350e7d4..57216983e4be 100644 --- a/gcc/config/gcn/gcn.opt +++ b/gcc/config/gcn/gcn.opt @@ -73,6 +73,12 @@ Target Report RejectNegative Joined UInteger Var(stack_size_opt) Init(-1) mlocal-symbol-id= Target RejectNegative Report JoinedOrMissing Var(local_symbol_id) Init(0) +int gang_local_size_opt = -1 + +mgang-local-size= +Target Report 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/doc/tm.texi b/gcc/doc/tm.texi index 055cda45862e..581b7b51eeb0 100644 --- a/gcc/doc/tm.texi +++ b/gcc/doc/tm.texi @@ -6214,7 +6214,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 @@ -6223,7 +6223,7 @@ 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. @@ -6234,6 +6234,13 @@ Define this hook to TRUE if arguments to offload regions should be exploded, i.e. passed as true arguments rather than in an argument array. @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 5dd84b2e3dc1..afa19d4ac63c 100644 --- a/gcc/doc/tm.texi.in +++ b/gcc/doc/tm.texi.in @@ -4225,6 +4225,8 @@ address; but often a machine-dependent strategy can generate better code. @hook TARGET_GOACC_EXPLODE_ARGS +@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 2b730d057781..35a19b23ff21 100644 --- a/gcc/omp-offload.c +++ b/gcc/omp-offload.c @@ -1631,7 +1631,7 @@ maybe_discard_oacc_function (tree decl) struct addr_expr_rewrite_info { gimple *stmt; - hash_set *adjusted_vars; + hash_map *adjusted_vars; bool avoid_pointer_conversion; bool modified; }; @@ -1645,19 +1645,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), @@ -1672,6 +1673,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; } @@ -1705,7 +1715,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; @@ -1861,8 +1872,104 @@ 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 (stmt); + 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; } @@ -1879,7 +1986,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 adjusted_vars; + hash_map adjusted_vars; /* Now lower internal loop functions to target-specific code sequences. */ @@ -1986,9 +2093,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; @@ -2163,22 +2272,9 @@ 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 && oacc_get_fn_attrib (fun->decl); }; virtual unsigned int execute (function *) diff --git a/gcc/omp-sese.c b/gcc/omp-sese.c index a64c9c3146a2..4dd3417066c6 100644 --- a/gcc/omp-sese.c +++ b/gcc/omp-sese.c @@ -54,6 +54,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. */ @@ -967,6 +969,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); } @@ -1004,7 +1008,8 @@ static void worker_single_copy (basic_block from, basic_block to, hash_set *def_escapes_block, hash_set *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 @@ -1015,10 +1020,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) @@ -1032,12 +1037,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, @@ -1206,13 +1222,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); @@ -1222,11 +1251,15 @@ worker_single_copy (basic_block from, basic_block to, gsi_insert_seq_after (&gsi, sender_seq, GSI_CONTINUE_LINKING); } +typedef hash_map > + blk_offset_map_t; + static void neuter_worker_single (parallel_g *par, unsigned outer_mask, bitmap worker_single, bitmap vector_single, vec *prop_set, - hash_set *partitioned_var_uses) + hash_set *partitioned_var_uses, + blk_offset_map_t *blk_offset_map) { unsigned mask = outer_mask | par->mask; @@ -1315,8 +1348,17 @@ 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); + { + auto 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); } @@ -1352,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 idx_decl_pair_t; + +typedef auto_vec 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; @@ -1449,8 +1634,120 @@ 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 > 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"); + auto 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); + auto 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 0b82a2417eb8..9aa25e8616ed 100644 --- a/gcc/omp-sese.h +++ b/gcc/omp-sese.h @@ -27,6 +27,6 @@ typedef auto_vec bb_pair_vec_t; extern void omp_find_sese (auto_vec &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 24932078732d..00421f3a6acd 100644 --- a/gcc/target.def +++ b/gcc/target.def @@ -1748,7 +1748,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 @@ -1756,7 +1756,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) DEFHOOK @@ -1771,6 +1771,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 dc4e80b4ae83..e9a81c58fb19 100644 --- a/gcc/targhooks.h +++ b/gcc/targhooks.h @@ -129,7 +129,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/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 000000000000..495fcde85119 --- /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" } */ + +#include +#include + +#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; +}