On Wed, 5 Dec 2018 21:10:45 +0000
Julian Brown <jul...@codesourcery.com> wrote:
> Thanks for review! How's this version?
>
> I took the liberty of fixing the patch for Fortran array-descriptor
> mappings that use a PSET, also, and adding another test for that
> functionality.
This is a ping/new version of this patch, incorporating previous review
comments and also fixing the inheritance behaviour for references (e.g.
for array slices of Fortran function arguments). I've also merged the
two patches sent below into one:
https://gcc.gnu.org/ml/gcc-patches/2018-08/msg01790.html
https://gcc.gnu.org/ml/gcc-patches/2018-08/msg01791.html
The second part having been conditionally approved based on approval of
the first already, and rebased.
Re-tested with offloading to NVPTX. OK?
Thanks,
Julian
2019-06-09 Julian Brown <jul...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
gcc/
* gimplify.c (oacc_array_mapping_info): New struct.
(gimplify_omp_ctx): Add decl_data_clause hash map.
(new_omp_context): Zero-initialise above.
(delete_omp_context): Delete above if allocated.
(gimplify_scan_omp_clauses): Scan for array mappings on data
constructs, and record in above map.
(gomp_oacc_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Handle data mappings (e.g.
array slices) declared in lexically-enclosing data constructs.
* omp-low.c (lower_omp_target): Allow decl for bias not to be
present in OpenACC context.
gcc/fortran/
* trans-openmp.c (gfc_omp_finish_clause): Don't raise error for
assumed-size array if present in a lexically-enclosing data
construct.
(gfc_omp_finish_clause): Guard addition of clauses for pointers
with DECL_P.
gcc/testsuite/
* c-c++-common/goacc/acc-data-chain.c: New test.
* gfortran.dg/goacc/pr70828.f90: New test.
* gfortran.dg/goacc/pr70828-2.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.
commit db3025e82a47ee8ca9fa8c87daa4eb9a7007fe75
Author: Julian Brown <jul...@codesourcery.com>
Date: Thu Aug 16 20:02:10 2018 -0700
Inheritance of array sections on data constructs
2018-08-28 Julian Brown <jul...@codesourcery.com>
Cesar Philippidis <ce...@codesourcery.com>
gcc/
* gimplify.c (oacc_array_mapping_info): New struct.
(gimplify_omp_ctx): Add decl_data_clause hash map.
(new_omp_context): Zero-initialise above.
(delete_omp_context): Delete above if allocated.
(gimplify_scan_omp_clauses): Scan for array mappings on data constructs, and record in above map.
(gomp_oacc_needs_data_present): New function.
(gimplify_adjust_omp_clauses_1): Handle data mappings (e.g. array
slices) declared in lexically-enclosing data constructs.
* omp-low.c (lower_omp_target): Allow decl for bias not to be present
in OpenACC context.
gcc/fortran/
* trans-openmp.c (gfc_omp_finish_clause): Don't raise error for
assumed-size array if present in a lexically-enclosing data construct.
(gfc_omp_finish_clause): Guard addition of clauses for pointers with
DECL_P.
gcc/testsuite/
* c-c++-common/goacc/acc-data-chain.c: New test.
* gfortran.dg/goacc/pr70828.f90: New test.
* gfortran.dg/goacc/pr70828-2.f90: New test.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr70828.c: New test.
* testsuite/libgomp.oacc-fortran/implicit_copy.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-2.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-3.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-4.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-5.f90: New test.
* testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 0eb5956cc53..56d56151a00 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1076,9 +1076,13 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
tree decl = OMP_CLAUSE_DECL (c);
- /* Assumed-size arrays can't be mapped implicitly, they have to be
- mapped explicitly using array sections. */
- if (TREE_CODE (decl) == PARM_DECL
+ /* Assumed-size arrays can't be mapped implicitly, they have to be mapped
+ explicitly using array sections. An exception is if the array is
+ mapped explicitly in an enclosing data construct for OpenACC, in which
+ case we see GOMP_MAP_FORCE_PRESENT here and do not need to raise an
+ error. */
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_PRESENT
+ && TREE_CODE (decl) == PARM_DECL
&& GFC_ARRAY_TYPE_P (TREE_TYPE (decl))
&& GFC_TYPE_ARRAY_AKIND (TREE_TYPE (decl)) == GFC_ARRAY_UNKNOWN
&& GFC_TYPE_ARRAY_UBOUND (TREE_TYPE (decl),
@@ -1091,7 +1095,7 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p)
}
tree c2 = NULL_TREE, c3 = NULL_TREE, c4 = NULL_TREE;
- if (POINTER_TYPE_P (TREE_TYPE (decl)))
+ if (DECL_P (decl) && POINTER_TYPE_P (TREE_TYPE (decl)))
{
if (!gfc_omp_privatize_by_reference (decl)
&& !GFC_DECL_GET_SCALAR_POINTER (decl)
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 0f3e917bc22..1b15b74550b 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -202,6 +202,18 @@ enum gimplify_defaultmap_kind
GDMK_POINTER
};
+/* Used to record clauses representing array slices on data directives that
+ may affect implicit mapping semantics on enclosed OpenACC parallel/kernels
+ regions. PSET is used for Fortran array slices with array descriptors,
+ or NULL otherwise. */
+struct oacc_array_mapping_info
+{
+ tree mapping;
+ tree pset;
+ tree pointer;
+ tree ref;
+};
+
struct gimplify_omp_ctx
{
struct gimplify_omp_ctx *outer_context;
@@ -217,6 +229,7 @@ struct gimplify_omp_ctx
bool distribute;
bool target_firstprivatize_array_bases;
int defaultmap[4];
+ hash_map<tree, oacc_array_mapping_info> *decl_data_clause;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -443,6 +456,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;
}
@@ -455,6 +469,8 @@ delete_omp_context (struct gimplify_omp_ctx *c)
splay_tree_delete (c->variables);
delete c->privatized_types;
c->loop_iter_var.release ();
+ if (c->decl_data_clause)
+ delete c->decl_data_clause;
XDELETE (c);
}
@@ -8421,8 +8437,57 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET:
break;
case OACC_DATA:
- if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
- break;
+ {
+ tree base_ptr = OMP_CLAUSE_CHAIN (c);
+ tree pset = NULL;
+ if (base_ptr
+ && OMP_CLAUSE_CODE (base_ptr) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (base_ptr) == GOMP_MAP_TO_PSET)
+ {
+ pset = base_ptr;
+ base_ptr = OMP_CLAUSE_CHAIN (base_ptr);
+ }
+ if (base_ptr
+ && OMP_CLAUSE_CODE (base_ptr) == OMP_CLAUSE_MAP
+ && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER))
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+ && ((OMP_CLAUSE_MAP_KIND (base_ptr)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ || OMP_CLAUSE_MAP_KIND (base_ptr) == GOMP_MAP_POINTER))
+ {
+ /* If we have an array descriptor, fish the right base
+ address variable to use out of that (otherwise we'd have
+ to deconstruct "arr.data" in the subsequent pointer
+ mapping). */
+ tree base_addr = pset ? OMP_CLAUSE_DECL (pset)
+ : OMP_CLAUSE_DECL (base_ptr);
+ if (!ctx->decl_data_clause)
+ ctx->decl_data_clause
+ = new hash_map<tree, oacc_array_mapping_info>;
+ oacc_array_mapping_info ai;
+ ai.mapping = unshare_expr (c);
+ ai.pset = pset ? unshare_expr (pset) : NULL;
+ ai.pointer = unshare_expr (base_ptr);
+ ai.ref = NULL_TREE;
+ if (TREE_CODE (base_addr) == INDIRECT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (base_addr, 0)))
+ == REFERENCE_TYPE))
+ {
+ base_addr = TREE_OPERAND (base_addr, 0);
+ tree ref_clause = OMP_CLAUSE_CHAIN (base_ptr);
+ gcc_assert ((OMP_CLAUSE_CODE (ref_clause)
+ == OMP_CLAUSE_MAP)
+ && (OMP_CLAUSE_MAP_KIND (ref_clause)
+ == GOMP_MAP_POINTER));
+ ai.ref = unshare_expr (ref_clause);
+ }
+ ctx->decl_data_clause->put (base_addr, ai);
+ }
+ if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+ break;
+ }
/* FALLTHRU */
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
@@ -9434,6 +9499,50 @@ struct gimplify_adjust_omp_clauses_data
gimple_seq *pre_p;
};
+/* For OpenACC parallel and kernels 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 parallel or
+ kernels region. This function returns a pointer to an
+ oacc_array_mapping_info if an array slice of DECL is specified in a
+ lexically-enclosing data construct, or returns NULL otherwise. */
+
+static oacc_array_mapping_info *
+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_KERNELS)
+ return NULL;
+
+ tree type = TREE_TYPE (decl);
+ if (TREE_CODE (type) == REFERENCE_TYPE)
+ type = TREE_TYPE (type);
+
+ if (TREE_CODE (type) != ARRAY_TYPE
+ && TREE_CODE (type) != POINTER_TYPE
+ && TREE_CODE (type) != RECORD_TYPE
+ && (TREE_CODE (type) != POINTER_TYPE
+ || TREE_CODE (TREE_TYPE (type)) != ARRAY_TYPE))
+ return NULL;
+
+ decl = get_base_address (decl);
+
+ for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context)
+ {
+ oacc_array_mapping_info *ret;
+
+ if (ctx->region_type != ORT_ACC_DATA)
+ break;
+
+ if (ctx->decl_data_clause && (ret = ctx->decl_data_clause->get (decl)))
+ return ret;
+ }
+
+ return NULL;
+}
+
/* For all variables that were not actually used within the context,
remove PRIVATE, SHARED, and FIRSTPRIVATE clauses. */
@@ -9534,6 +9643,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;
+ oacc_array_mapping_info *array_info;
if (private_debug)
OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
@@ -9542,6 +9652,69 @@ 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 ((code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE)
+ && (array_info = gomp_oacc_needs_data_present (decl)))
+ {
+ tree mapping = array_info->mapping;
+ tree pointer = array_info->pointer;
+
+ 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);
+
+ OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping));
+ OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
+ OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping));
+
+ /* Create a new data clause for the firstprivate pointer. */
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (pointer));
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_POINTER);
+
+ /* For GOMP_MAP_FIRSTPRIVATE_POINTER, this is a bias, not a size. */
+ OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (pointer));
+
+ /* Create a new data clause for the PSET, if present. */
+ tree psetc = NULL;
+ if (array_info->pset)
+ {
+ tree pset = array_info->pset;
+ psetc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (psetc) = unshare_expr (OMP_CLAUSE_DECL (pset));
+ OMP_CLAUSE_SIZE (psetc) = unshare_expr (OMP_CLAUSE_SIZE (pset));
+ OMP_CLAUSE_SET_MAP_KIND (psetc, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_CHAIN (psetc) = nc;
+ }
+
+ gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ gimplify_expr (&OMP_CLAUSE_DECL (clause), pre_p, NULL,
+ is_gimple_lvalue, fb_lvalue);
+ gimplify_expr (&OMP_CLAUSE_SIZE (clause), pre_p, NULL,
+ is_gimple_val, fb_rvalue);
+ gimplify_expr (&OMP_CLAUSE_SIZE (nc), pre_p, NULL, is_gimple_val,
+ fb_rvalue);
+ gimplify_omp_ctxp = ctx;
+
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
+ OMP_CLAUSE_CHAIN (clause) = psetc ? psetc : nc;
+
+ if (array_info->ref)
+ {
+ tree refc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (refc)
+ = unshare_expr (OMP_CLAUSE_DECL (array_info->ref));
+ OMP_CLAUSE_SIZE (refc)
+ = unshare_expr (OMP_CLAUSE_SIZE (array_info->ref));
+ OMP_CLAUSE_SET_MAP_KIND (refc, GOMP_MAP_POINTER);
+ OMP_CLAUSE_CHAIN (refc) = OMP_CLAUSE_CHAIN (nc);
+ OMP_CLAUSE_CHAIN (nc) = refc;
+ }
+ }
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)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index a7f35ffe416..9c0042bd8c2 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -10445,8 +10445,14 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
x = fold_convert_loc (clause_loc, type, x);
if (!integer_zerop (OMP_CLAUSE_SIZE (c)))
{
- tree bias = OMP_CLAUSE_SIZE (c);
- if (DECL_P (bias))
+ tree bias = OMP_CLAUSE_SIZE (c), remapped_bias;
+ if (is_gimple_omp_oacc (ctx->stmt))
+ {
+ if (DECL_P (bias)
+ && (remapped_bias = maybe_lookup_decl (bias, ctx)))
+ bias = remapped_bias;
+ }
+ else if (DECL_P (bias))
bias = lookup_decl (bias, ctx);
bias = fold_convert_loc (clause_loc, sizetype, bias);
bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype,
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 00000000000..8a039be15f5
--- /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.alloc:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(alloc:a \\\[pointer assign, bias: 0\\\]\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90
new file mode 100644
index 00000000000..2e581205c8a
--- /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:MEM\\\[\\(c_char \\*\\)\_\[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:MEM\\\[\\(c_char \\*\\)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 00000000000..357114ccfd3
--- /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 00000000000..4b6dbd7538f
--- /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/implicit_copy.f90 b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90
new file mode 100644
index 00000000000..7a99f294c85
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90
@@ -0,0 +1,30 @@
+! { dg-do run }
+
+integer function test()
+ implicit none
+ integer, parameter :: n = 10
+ real(8), dimension(n) :: a, b, c
+ integer i
+
+ do i = 1, n
+ a(i) = i
+ b(i) = 1
+ end do
+
+ !$acc data copyin(a(1:n), b(1:n))
+ !$acc parallel loop
+ do i = 1, n
+ c(i) = a(i) * b(i)
+ end do
+ !$acc end data
+
+ do i = 1, n
+ if (c(i) /= a(i) * b(i)) call abort
+ end do
+end function test
+
+program main
+ implicit none
+ integer i, test
+ i = test()
+end program main
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 00000000000..22a956622bb
--- /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 00000000000..ff17d10cfa3
--- /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 00000000000..01da999b33d
--- /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 00000000000..8a16e3d5550
--- /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 00000000000..e99c3649159
--- /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 00000000000..f87d232fe42
--- /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
