public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-13] OpenACC: Reimplement "inheritance" for lexically-nested offload regions
@ 2023-06-19 22:17 Julian Brown
0 siblings, 0 replies; only message in thread
From: Julian Brown @ 2023-06-19 22:17 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:7606557e05f1f9578295f4a7c22f309b2b9229b9
commit 7606557e05f1f9578295f4a7c22f309b2b9229b9
Author: Julian Brown <julian@codesourcery.com>
Date: Thu Jun 15 19:23:10 2023 +0000
OpenACC: Reimplement "inheritance" for lexically-nested offload regions
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 <julian@codesourcery.com>
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.
Diff:
---
gcc/ChangeLog.omp | 17 ++
gcc/c-family/ChangeLog.omp | 5 +
gcc/c-family/c-omp.cc | 13 +-
gcc/gimplify.cc | 208 ++++++++++++++++-----
gcc/testsuite/ChangeLog.omp | 6 +
gcc/testsuite/c-c++-common/goacc/acc-data-chain.c | 4 +-
gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 | 35 ++++
gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 3 +-
libgomp/ChangeLog.omp | 11 ++
.../libgomp.oacc-c-c++-common/pr70828-2.c | 2 -
.../testsuite/libgomp.oacc-c-c++-common/pr70828.c | 2 -
.../testsuite/libgomp.oacc-fortran/pr70828-2.f90 | 2 -
.../testsuite/libgomp.oacc-fortran/pr70828-3.f90 | 2 -
.../testsuite/libgomp.oacc-fortran/pr70828-4.f90 | 2 -
.../testsuite/libgomp.oacc-fortran/pr70828-5.f90 | 2 -
.../testsuite/libgomp.oacc-fortran/pr70828-6.f90 | 2 -
libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 | 2 -
17 files changed, 241 insertions(+), 77 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index f4dfc6cc6cd..32f86dd5c5f 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,20 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* gimplify.cc (omp_segregate_mapping_groups): Handle "present" groups.
diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 32f8febbb11..ebafb8f17df 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,8 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit
+ pointer nodes for OpenACC.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* c-common.h (c_omp_region_type): Add C_ORT_ACC_TARGET.
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<tree, omp_mapping_group *> *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<omp_mapping_group> *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<tree, omp_mapping_group *>;
+
+ 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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index 0d89832a180..e9661d9b3cf 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,9 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* gfortran.dg/gomp/map-12.f90: Adjust scan output.
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/ChangeLog.omp b/libgomp/ChangeLog.omp
index a7676907d8d..8817c7a500f 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,14 @@
+2023-06-19 Julian Brown <julian@codesourcery.com>
+
+ * 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.
+
2023-06-19 Julian Brown <julian@codesourcery.com>
* oacc-mem.c (find_group_last, goacc_enter_data_internal,
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 } }
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2023-06-19 22:17 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-06-19 22:17 [gcc/devel/omp/gcc-13] OpenACC: Reimplement "inheritance" for lexically-nested offload regions Julian Brown
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).