https://gcc.gnu.org/g:f170e172ab2f559aa765a28af796c1a4897adb00
commit f170e172ab2f559aa765a28af796c1a4897adb00
Author: Julian Brown <jul...@codesourcery.com>
Date: Sat Apr 19 23:41:41 2025 +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.
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: New test.
* gfortran.dg/goacc/pr70828.f90: Likewise.
* gfortran.dg/goacc/assumed-size.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: Likewise.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: Likewise.
Diff:
---
gcc/c-family/c-omp.cc | 13 +-
gcc/gimplify.cc | 207 ++++++++++++++++-----
gcc/testsuite/c-c++-common/goacc/acc-data-chain.c | 24 +++
gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 | 35 ++++
gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 22 +++
.../libgomp.oacc-c-c++-common/pr70828-2.c | 34 ++++
.../testsuite/libgomp.oacc-c-c++-common/pr70828.c | 27 +++
.../testsuite/libgomp.oacc-fortran/pr70828-2.f90 | 31 +++
.../testsuite/libgomp.oacc-fortran/pr70828-3.f90 | 34 ++++
.../testsuite/libgomp.oacc-fortran/pr70828-4.f90 | 31 +++
.../testsuite/libgomp.oacc-fortran/pr70828-5.f90 | 29 +++
.../testsuite/libgomp.oacc-fortran/pr70828-6.f90 | 28 +++
libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 | 24 +++
13 files changed, 484 insertions(+), 55 deletions(-)
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 65cb3e2ccb4f..25fadf1bb191 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -3868,7 +3868,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_p
+ if (openmp_p
+ && !target_p
&& (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr))
== ARRAY_TYPE))
break;
@@ -3879,7 +3880,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_p)
+ if (decl_p && (!openmp_p || target_p))
{
/* See comment for ACCESS_INDEXED_REF_TO_ARRAY above. */
enum gomp_map_kind k = chain_p ? GOMP_MAP_POINTER
@@ -3935,9 +3936,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_p || !openmp_p) && !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 23a1cbf48746..31e4b4d97b0c 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -272,6 +272,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;
@@ -294,6 +336,7 @@ struct gimplify_omp_ctx
bool in_call_args;
bool ompacc;
int defaultmap[5];
+ hash_map<tree, omp_mapping_group *> *decl_data_clause;
};
struct privatize_reduction
@@ -528,6 +571,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;
}
@@ -540,6 +584,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);
}
@@ -10310,18 +10355,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. */
@@ -10338,26 +10371,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)
{
@@ -10619,16 +10632,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;
@@ -13646,6 +13650,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
&& TREE_CODE (TREE_TYPE (basetype)) == POINTER_TYPE)
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
@@ -14385,11 +14401,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
@@ -14530,6 +14541,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. */
@@ -14651,6 +14708,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))
@@ -14659,6 +14717,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)
@@ -15132,9 +15242,6 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p,
gimple_seq body, tree *list_p,
switch (code)
{
case OACC_DATA:
- if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
- break;
- /* Fallthrough. */
case OACC_HOST_DATA:
case OACC_ENTER_DATA:
case OACC_EXIT_DATA:
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
new file mode 100644
index 000000000000..622f1992f88c
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
@@ -0,0 +1,24 @@
+/* Ensure that the gimplifier does not remove any existing clauses as
+ it inserts new implicit data clauses. */
+
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+static int a[N], b[N];
+
+int main(int argc, char *argv[])
+{
+ int i;
+
+#pragma acc data copyin(a[0:N]) copyout (b[0:N])
+ {
+#pragma acc parallel loop
+ for (i = 0; i < N; i++)
+ b[i] = a[i];
+ }
+
+ return 0;
+}
+
+// { dg-final { scan-tree-dump-times "omp target oacc_data
map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len:
400\\\]\\)" 1 "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 000000000000..4fced2e70c9e
--- /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
new file mode 100644
index 000000000000..fcfe0865fc41
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90
@@ -0,0 +1,22 @@
+! Ensure that pointer mappings are preserved in nested parallel
+! constructs.
+
+! { dg-additional-options "-fdump-tree-gimple" }
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ !$acc data copy(data(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ data(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+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" } }
+! { 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
new file mode 100644
index 000000000000..357114ccfd38
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c
@@ -0,0 +1,34 @@
+/* Subarray declared on data construct, accessed through pointer. */
+
+#include <assert.h>
+
+void
+s1 (int *arr, int c)
+{
+#pragma acc data copy(arr[5:c-10])
+ {
+#pragma acc parallel loop
+ for (int i = 5; i < c - 5; i++)
+ arr[i] = i;
+ }
+}
+
+int
+main (int argc, char* argv[])
+{
+ const int c = 100;
+ int arr[c];
+
+ for (int i = 0; i < c; i++)
+ arr[i] = 0;
+
+ s1 (arr, c);
+
+ for (int i = 0; i < c; i++)
+ if (i >= 5 && i < c - 5)
+ assert (arr[i] == i);
+ else
+ assert (arr[i] == 0);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
new file mode 100644
index 000000000000..4b6dbd7538fb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
@@ -0,0 +1,27 @@
+/* Subarray declared on enclosing data construct. */
+
+#include <assert.h>
+
+int
+main ()
+{
+ int a[100], i;
+
+ for (i = 0; i < 100; i++)
+ a[i] = 0;
+
+#pragma acc data copy(a[10:80])
+ {
+ #pragma acc parallel loop
+ for (i = 10; i < 90; i++)
+ a[i] = i;
+ }
+
+ for (i = 0; i < 100; i++)
+ if (i >= 10 && i < 90)
+ assert (a[i] == i);
+ else
+ assert (a[i] == 0);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
new file mode 100644
index 000000000000..22a956622bb2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
@@ -0,0 +1,31 @@
+! Subarrays declared on data construct: assumed-shape array.
+
+subroutine s1(n, arr)
+ integer :: n
+ integer :: arr(n)
+
+ !$acc data copy(arr(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ arr(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+end subroutine s1
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ call s1(n, data)
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
new file mode 100644
index 000000000000..ff17d10cfa31
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
@@ -0,0 +1,34 @@
+! Subarrays declared on data construct: deferred-shape array.
+
+subroutine s1(n, arr)
+ integer :: n
+ integer :: arr(n)
+
+ !$acc data copy(arr(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ arr(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+end subroutine s1
+
+program test
+ integer, parameter :: n = 100
+ integer i
+ integer, allocatable :: data(:)
+
+ allocate (data(1:n))
+
+ data(:) = 0
+
+ call s1(n, data)
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
new file mode 100644
index 000000000000..01da999b33d7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
@@ -0,0 +1,31 @@
+! Subarrays declared on data construct: assumed-size array.
+
+subroutine s1(n, arr)
+ integer :: n
+ integer :: arr(*)
+
+ !$acc data copy(arr(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ arr(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+end subroutine s1
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ call s1(n, data)
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
new file mode 100644
index 000000000000..8a16e3d5550a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
@@ -0,0 +1,29 @@
+! Subarrays on parallel construct (no data construct): assumed-size array.
+
+subroutine s1(n, arr)
+ integer :: n
+ integer :: arr(*)
+
+ !$acc parallel loop copy(arr(5:n-10))
+ do i = 10, n - 10
+ arr(i) = i
+ end do
+ !$acc end parallel loop
+end subroutine s1
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ call s1(n, data)
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
new file mode 100644
index 000000000000..e99c36491597
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
@@ -0,0 +1,28 @@
+! Subarrays declared on data construct: allocatable array (with array
+! descriptor).
+
+program test
+ integer, parameter :: n = 100
+ integer i
+ integer, allocatable :: data(:)
+
+ allocate (data(1:n))
+
+ data(:) = 0
+
+ !$acc data copy(data(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ data(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
new file mode 100644
index 000000000000..f87d232fe42c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
@@ -0,0 +1,24 @@
+! Subarrays on data construct: explicit-shape array.
+
+program test
+ integer, parameter :: n = 100
+ integer i, data(n)
+
+ data(:) = 0
+
+ !$acc data copy(data(5:n-10))
+ !$acc parallel loop
+ do i = 10, n - 10
+ data(i) = i
+ end do
+ !$acc end parallel loop
+ !$acc end data
+
+ do i = 1, n
+ if ((i < 10 .or. i > n-10)) then
+ if ((data(i) .ne. 0)) call abort
+ else if (data(i) .ne. i) then
+ call abort
+ end if
+ end do
+end program test
