From patchwork Mon Jun 19 21:17:35 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Julian Brown X-Patchwork-Id: 1796818 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@legolas.ozlabs.org Authentication-Results: legolas.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+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) 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 ECDSA (P-384) server-digest SHA384) (No client certificate requested) by legolas.ozlabs.org (Postfix) with ESMTPS id 4QlN5x5XhNz20XS for ; Tue, 20 Jun 2023 07:23:01 +1000 (AEST) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id BFA323888C45 for ; Mon, 19 Jun 2023 21:22:59 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id BAB753882168; Mon, 19 Jun 2023 21:19:41 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org BAB753882168 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.00,255,1681200000"; d="scan'208";a="9264461" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa4.mentor.iphmx.com with ESMTP; 19 Jun 2023 13:19:40 -0800 IronPort-SDR: IJlTwJpSPkZNqy81U7kJQymp+FmxWN1iDliRr7Hjq4GA+eDmKeOesVuV5nuR+vv7xxfhuRKZnp Qb003T34xBMwZeE9YRuxkEYTVbkWwRXt2i+ikqjtrZNY/sdEYvFL6MFCLI8nhw27v+FTCwBaI4 n51HZJpUxKvcsrQQkGmvYM/nOI3IoZs8blAsnEVXS4iqYic166IbmGdvYq9CBJu+T/sLynC93l Ml0XRhYJva/11H50LMkWh4NquF0AlFUfKVlRfpOgztjKLpRJQn5Of044nHwaCwjCws24F4Z6dG kLs= From: Julian Brown To: CC: , , Subject: [PATCH 11/14] OpenACC: Reimplement "inheritance" for lexically-nested offload regions Date: Mon, 19 Jun 2023 21:17:35 +0000 Message-ID: <77c0972c0e363b485b5fe1aa57f40221794a25ac.1687201316.git.julian@codesourcery.com> X-Mailer: git-send-email 2.25.1 In-Reply-To: References: MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) 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: , Errors-To: gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" This patch reimplements "lexical inheritance" for OpenACC offload regions inside "data" regions, allowing e.g. this to work: int *ptr; [...] #pragma acc data copyin(ptr[10:2]) { #pragma acc parallel { ... } } here, the "copyin" is mirrored on the inner "acc parallel" as "present(ptr[10:2])" -- allowing code within the parallel to use that section of the array even though the mapping is implicit. In terms of implementation, this works by expanding mapping nodes for "acc data" to include pointer mappings that might be needed by inner offload regions. The resulting mapping group is then copied to the inner offload region as needed, rewriting the first node to "force_present". The pointer mapping nodes are then removed from the "acc data" later during gimplification. For OpenMP, pointer mapping nodes on equivalent "omp data" regions are not needed, so remain suppressed during expansion. 2023-06-16 Julian Brown gcc/c-family/ * c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit pointer nodes for OpenACC. gcc/ * gimplify.cc (omp_tsort_mark, omp_mapping_group): Move before gimplify_omp_ctx. Add constructor to omp_mapping_group. (gimplify_omp_ctx): Add DECL_DATA_CLAUSE field. (new_omp_context, delete_omp_context): Initialise and free above field. (omp_gather_mapping_groups_1): Use constructor for omp_mapping_group. (gimplify_scan_omp_clauses): Record mappings that might be lexically inherited. Don't remove GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE yet. (gomp_oacc_needs_data_present): New function. (gimplify_adjust_omp_clauses_1): Implement lexical inheritance behaviour for OpenACC. (gimplify_adjust_omp_clauses): Remove GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE here instead, after lexical inheritance is done. gcc/testsuite/ * c-c++-common/goacc/acc-data-chain.c: Re-enable scan test. * gfortran.dg/goacc/pr70828.f90: Likewise. * gfortran.dg/goacc/assumed-size.f90: New test. libgomp/ * testsuite/libgomp.oacc-c-c++-common/pr70828.c: Un-XFAIL. * testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-2.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-3.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-4.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-5.f90: Un-XFAIL. * testsuite/libgomp.oacc-fortran/pr70828-6.f90: Un-XFAIL. --- gcc/c-family/c-omp.cc | 13 +- gcc/gimplify.cc | 208 +++++++++++++----- .../c-c++-common/goacc/acc-data-chain.c | 4 +- .../gfortran.dg/goacc/assumed-size.f90 | 35 +++ gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 3 +- .../libgomp.oacc-c-c++-common/pr70828-2.c | 2 - .../libgomp.oacc-c-c++-common/pr70828.c | 2 - .../libgomp.oacc-fortran/pr70828-2.f90 | 2 - .../libgomp.oacc-fortran/pr70828-3.f90 | 2 - .../libgomp.oacc-fortran/pr70828-4.f90 | 2 - .../libgomp.oacc-fortran/pr70828-5.f90 | 2 - .../libgomp.oacc-fortran/pr70828-6.f90 | 2 - .../libgomp.oacc-fortran/pr70828.f90 | 2 - 13 files changed, 202 insertions(+), 77 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index e55b2aec920..291a26293ef 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4313,7 +4313,8 @@ c_omp_address_inspector::expand_array_base (tree c, /* The code handling "firstprivatize_array_bases" in gimplify.cc is relevant here. What do we need to create for arrays at this stage? (This condition doesn't feel quite right. FIXME?) */ - if (!target + if (openmp + && !target && (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr)) == ARRAY_TYPE)) break; @@ -4324,7 +4325,7 @@ c_omp_address_inspector::expand_array_base (tree c, virtual_origin); tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); - if (decl_p && target) + if (decl_p && (!openmp || target)) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); else { @@ -4375,9 +4376,11 @@ c_omp_address_inspector::expand_array_base (tree c, tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr); c2 = build_omp_clause (loc, OMP_CLAUSE_MAP); /* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute - regions (e.g. "acc data" constructs). It'll be removed anyway in - gimplify.cc, but doing it this way maintains diagnostic - behaviour. */ + regions (e.g. "acc data" constructs). It is used during "lexical + inheritance" of mapping clauses on enclosed target + (parallel/serial/kernels) regions, i.e. creating "present" mappings + for sections of pointer-based arrays. It's also used for + diagnostics. */ if (decl_p && (target || !openmp) && !chain_p && !declare_target_p) OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER); else diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index e21e9d99cc9..55befa4d3c1 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -218,6 +218,48 @@ enum gimplify_defaultmap_kind GDMK_POINTER }; +/* Used for topological sorting of mapping groups. UNVISITED means we haven't + started processing the group yet. The TEMPORARY mark is used when we first + encounter a group on a depth-first traversal, and the PERMANENT mark is used + when we have processed all the group's children (i.e. all the base pointers + referred to by the group's mapping nodes, recursively). */ + +enum omp_tsort_mark { + UNVISITED, + TEMPORARY, + PERMANENT +}; + +/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map" + clause. */ + +struct omp_mapping_group { + tree *grp_start; + tree grp_end; + omp_tsort_mark mark; + /* If we've removed the group but need to reindex, mark the group as + deleted. */ + bool deleted; + /* The group points to an already-created "GOMP_MAP_STRUCT + GOMP_MAP_ATTACH_DETACH" pair. */ + bool reprocess_struct; + /* The group should use "zero-length" allocations for pointers that are not + mapped "to" on the same directive. */ + bool fragile; + struct omp_mapping_group *sibling; + struct omp_mapping_group *next; + + omp_mapping_group (tree *_start, tree _end) + : grp_start (_start), grp_end (_end), mark (UNVISITED), deleted (false), + reprocess_struct (false), fragile (false), sibling (NULL), next (NULL) + { + } + + omp_mapping_group () + { + } +}; + struct gimplify_omp_ctx { struct gimplify_omp_ctx *outer_context; @@ -239,6 +281,7 @@ struct gimplify_omp_ctx bool in_for_exprs; bool ompacc; int defaultmap[5]; + hash_map *decl_data_clause; }; struct privatize_reduction @@ -473,6 +516,7 @@ new_omp_context (enum omp_region_type region_type) c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP; c->defaultmap[GDMK_ALLOCATABLE] = GOVD_MAP; c->defaultmap[GDMK_POINTER] = GOVD_MAP; + c->decl_data_clause = NULL; return c; } @@ -485,6 +529,7 @@ delete_omp_context (struct gimplify_omp_ctx *c) splay_tree_delete (c->variables); delete c->privatized_types; c->loop_iter_var.release (); + delete c->decl_data_clause; XDELETE (c); } @@ -8988,18 +9033,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp, return base; } -/* Used for topological sorting of mapping groups. UNVISITED means we haven't - started processing the group yet. The TEMPORARY mark is used when we first - encounter a group on a depth-first traversal, and the PERMANENT mark is used - when we have processed all the group's children (i.e. all the base pointers - referred to by the group's mapping nodes, recursively). */ - -enum omp_tsort_mark { - UNVISITED, - TEMPORARY, - PERMANENT -}; - /* Hash for trees based on operand_equal_p. Like tree_operand_hash but ignores side effects in the equality comparisons. */ @@ -9016,26 +9049,6 @@ tree_operand_hash_no_se::equal (const value_type &t1, return operand_equal_p (t1, t2, OEP_MATCH_SIDE_EFFECTS); } -/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map" - clause. */ - -struct omp_mapping_group { - tree *grp_start; - tree grp_end; - omp_tsort_mark mark; - /* If we've removed the group but need to reindex, mark the group as - deleted. */ - bool deleted; - /* The group points to an already-created "GOMP_MAP_STRUCT - GOMP_MAP_ATTACH_DETACH" pair. */ - bool reprocess_struct; - /* The group should use "zero-length" allocations for pointers that are not - mapped "to" on the same directive. */ - bool fragile; - struct omp_mapping_group *sibling; - struct omp_mapping_group *next; -}; - DEBUG_FUNCTION void debug_mapping_group (omp_mapping_group *grp) { @@ -9276,16 +9289,7 @@ omp_gather_mapping_groups_1 (tree *list_p, vec *groups, continue; tree *grp_last_p = omp_group_last (cp); - omp_mapping_group grp; - - grp.grp_start = cp; - grp.grp_end = *grp_last_p; - grp.mark = UNVISITED; - grp.sibling = NULL; - grp.deleted = false; - grp.reprocess_struct = false; - grp.fragile = false; - grp.next = NULL; + omp_mapping_group grp (cp, *grp_last_p); groups->safe_push (grp); cp = grp_last_p; @@ -12267,6 +12271,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, } break; } + if (code == OACC_DATA && *grp_start_p != grp_end) + { + if (!ctx->decl_data_clause) + ctx->decl_data_clause = new hash_map; + + omp_mapping_group *grp + = new omp_mapping_group (grp_start_p, grp_end); + + gcc_assert (DECL_P (decl)); + + ctx->decl_data_clause->put (decl, grp); + } flags = GOVD_MAP | GOVD_EXPLICIT; if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM) @@ -12953,11 +12969,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, gcc_unreachable (); } - if (code == OACC_DATA - && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP - && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - remove = true; if (remove) *list_p = OMP_CLAUSE_CHAIN (c); else @@ -13098,6 +13109,52 @@ struct gimplify_adjust_omp_clauses_data gimple_seq *pre_p; }; +/* For OpenACC offload regions, the implicit data mappings for arrays must + respect explicit data clauses set by a containing acc data region. + Specifically, an array section on the data clause must be transformed into + an equivalent PRESENT mapping on the inner offload region. + This function returns a pointer to a mapping group if an array slice of DECL + is specified on a lexically-enclosing data construct, or returns NULL + otherwise. */ + +static omp_mapping_group * +gomp_oacc_needs_data_present (tree decl) +{ + gimplify_omp_ctx *ctx = NULL; + + if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL + && gimplify_omp_ctxp->region_type != ORT_ACC_SERIAL + && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS) + return NULL; + + if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE + && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE + && TREE_CODE (TREE_TYPE (decl)) != RECORD_TYPE + && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE + || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE)) + return NULL; + + decl = get_base_address (decl); + + for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context) + { + splay_tree_node on + = splay_tree_lookup (ctx->variables, (splay_tree_key) decl); + + if (ctx->region_type == ORT_ACC_DATA + && on + && (((int) on->value) & GOVD_EXPLICIT) + && ctx->decl_data_clause != NULL) + { + omp_mapping_group **pgrp = ctx->decl_data_clause->get (decl); + if (pgrp) + return *pgrp; + } + } + + return NULL; +} + /* For all variables that were not actually used within the context, remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */ @@ -13219,6 +13276,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) clause = build_omp_clause (input_location, code); OMP_CLAUSE_DECL (clause) = decl; OMP_CLAUSE_CHAIN (clause) = chain; + omp_mapping_group *outer_grp; if (private_debug) OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1; else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF)) @@ -13227,6 +13285,58 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) && (flags & GOVD_WRITTEN) == 0 && omp_shared_to_firstprivate_optimizable_decl_p (decl)) OMP_CLAUSE_SHARED_READONLY (clause) = 1; + else if ((gimplify_omp_ctxp->region_type & ORT_ACC) != 0 + && (code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE) + && (outer_grp = gomp_oacc_needs_data_present (decl))) + { + if (code == OMP_CLAUSE_FIRSTPRIVATE) + /* Oops, we have the wrong type of clause. Rebuild it. */ + clause = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + + tree mapping = *outer_grp->grp_start; + + OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT); + OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping)); + OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping)); + + /* Copy subsequent nodes (that are part of the mapping group) after the + initial one from the outer "acc data" directive -- "pointer" nodes, + including firstprivate_reference, pointer sets, etc. */ + + tree ptr = OMP_CLAUSE_CHAIN (mapping); + tree *ins = &OMP_CLAUSE_CHAIN (clause); + tree sentinel = OMP_CLAUSE_CHAIN (outer_grp->grp_end); + for (; ptr && ptr != sentinel; ptr = OMP_CLAUSE_CHAIN (ptr)) + { + tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause), + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (nc, OMP_CLAUSE_MAP_KIND (ptr)); + OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (ptr)); + OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (ptr)); + *ins = nc; + ins = &OMP_CLAUSE_CHAIN (nc); + } + + *ins = chain; + + gimplify_omp_ctx *ctx = gimplify_omp_ctxp; + gimplify_omp_ctxp = ctx->outer_context; + for (ptr = clause; ptr != chain; ptr = OMP_CLAUSE_CHAIN (ptr)) + { + /* The condition is specifically to not gimplify here if we have a + DECL_P with a DECL_VALUE_EXPR -- i.e. a VLA, or variable-sized + array section. If we do, omp-low.cc does not see the DECL_P it + expects here for e.g. firstprivate_pointer or + firstprivate_reference. */ + if (!DECL_P (OMP_CLAUSE_DECL (ptr))) + gimplify_expr (&OMP_CLAUSE_DECL (ptr), pre_p, NULL, + is_gimple_lvalue, fb_lvalue); + gimplify_expr (&OMP_CLAUSE_SIZE (ptr), pre_p, NULL, + is_gimple_val, fb_rvalue); + } + gimplify_omp_ctxp = ctx; + } else if (code == OMP_CLAUSE_FIRSTPRIVATE && (flags & GOVD_EXPLICIT) == 0) OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1; else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0) @@ -13689,16 +13799,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, case OMP_TARGET: break; case OACC_DATA: - if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE) - break; - goto check_firstprivate; case OACC_ENTER_DATA: case OACC_EXIT_DATA: case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: case OACC_HOST_DATA: - check_firstprivate: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c index 932786cec76..622f1992f88 100644 --- a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c +++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c @@ -21,6 +21,4 @@ int main(int argc, char *argv[]) } // { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len: 400\\\]\\)" 1 "gimple" } } -/* This isn't expected to work while the "lexical inheritance" support is - reverted. */ -// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map.alloc:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(alloc:a \\\[pointer assign, bias: 0\\\]\\)" 0 "gimple" } } +// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:a \\\[pointer assign, bias: 0\\\]\\)" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 b/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 new file mode 100644 index 00000000000..4fced2e70c9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 @@ -0,0 +1,35 @@ +! Test if implicitly determined data clauses work with an +! assumed-sized array variable. Note that the array variable, 'a', +! has been explicitly copied in and out via acc enter data and acc +! exit data, respectively. + +! This does not appear to be supported by the OpenACC standard as of version +! 3.0. Check for an appropriate error message. + +program test + implicit none + + integer, parameter :: n = 100 + integer a(n), i + + call dtest (a, n) + + do i = 1, n + if (a(i) /= i) call abort + end do +end program test + +subroutine dtest (a, n) + integer i, n + integer a(*) + + !$acc enter data copyin(a(1:n)) + + !$acc parallel loop +! { dg-error {implicit mapping of assumed size array 'a'} "" { target *-*-* } .-1 } + do i = 1, n + a(i) = i + end do + + !$acc exit data copyout(a(1:n)) +end subroutine dtest diff --git a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 index 72b0d9ae92c..fcfe0865fc4 100644 --- a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 @@ -19,5 +19,4 @@ program test end program test ! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(tofrom:data\\\[\_\[0-9\]+\\\] \\\[len: _\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: _\[0-9\]+\\\]\\)" 1 "gimple" } } -! Disable for now -! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 0 "gimple" } } +! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 1 "gimple" } } diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c index da5bb3f93c3..357114ccfd3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c @@ -32,5 +32,3 @@ main (int argc, char* argv[]) return 0; } - -/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c index 85d09bff1df..4b6dbd7538f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c @@ -25,5 +25,3 @@ main () return 0; } - -/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */ diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 index 2892b3d5938..22a956622bb 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 @@ -29,5 +29,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 index e28193b1a22..ff17d10cfa3 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 @@ -32,5 +32,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90 index 918295d5c8b..01da999b33d 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90 @@ -29,5 +29,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 index 3b5d05d1379..8a16e3d5550 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 @@ -27,5 +27,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 index d48168b22eb..e99c3649159 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 @@ -26,5 +26,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 index 5db49e1a569..f87d232fe42 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 @@ -22,5 +22,3 @@ program test end if end do end program test - -! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }