public inbox for gcc-cvs@sourceware.org
help / color / mirror / Atom feed
* [gcc/devel/omp/gcc-11] Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828)
@ 2021-05-13 16:09 Kwok Yeung
0 siblings, 0 replies; only message in thread
From: Kwok Yeung @ 2021-05-13 16:09 UTC (permalink / raw)
To: gcc-cvs
https://gcc.gnu.org/g:8b849547b9e10807d2ad75c4a7d91ce3cf4ff531
commit 8b849547b9e10807d2ad75c4a7d91ce3cf4ff531
Author: Julian Brown <julian@codesourcery.com>
Date: Tue Feb 26 14:22:41 2019 -0800
Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828)
2018-08-28 Julian Brown <julian@codesourcery.com>
Cesar Philippidis <cesar@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/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:
---
gcc/ChangeLog.omp | 15 +++
gcc/fortran/trans-openmp.c | 10 +-
gcc/gimplify.c | 143 ++++++++++++++++++++-
gcc/omp-low.c | 10 +-
gcc/testsuite/ChangeLog.omp | 7 +
gcc/testsuite/c-c++-common/goacc/acc-data-chain.c | 24 ++++
gcc/testsuite/gfortran.dg/goacc/pr70828.f90 | 22 ++++
libgomp/ChangeLog.omp | 12 ++
.../libgomp.oacc-c-c++-common/pr70828-2.c | 34 +++++
.../testsuite/libgomp.oacc-c-c++-common/pr70828.c | 27 ++++
.../libgomp.oacc-fortran/implicit_copy.f90 | 30 +++++
.../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 ++++
17 files changed, 504 insertions(+), 7 deletions(-)
diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp
index 5c8f877e18b..3cb13e99ad1 100644
--- a/gcc/ChangeLog.omp
+++ b/gcc/ChangeLog.omp
@@ -1,3 +1,18 @@
+2018-08-28 Julian Brown <julian@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-10-05 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index c6b8a22b737..1a99a453b12 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -1293,9 +1293,13 @@ gfc_omp_finish_clause (tree c, gimple_seq *pre_p, bool openacc)
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),
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index df6fdcc2c3e..24e4af364d5 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -217,6 +217,17 @@ 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;
+};
+
struct gimplify_omp_ctx
{
struct gimplify_omp_ctx *outer_context;
@@ -237,6 +248,7 @@ struct gimplify_omp_ctx
bool has_depend;
bool in_for_exprs;
int defaultmap[4];
+ hash_map<tree, oacc_array_mapping_info> *decl_data_clause;
};
static struct gimplify_ctx *gimplify_ctxp;
@@ -464,6 +476,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;
}
@@ -476,6 +489,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);
}
@@ -8959,8 +8974,41 @@ 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_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);
+ 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:
@@ -10197,6 +10245,46 @@ 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;
+
+ 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)
+ {
+ 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. */
@@ -10320,6 +10408,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))
@@ -10328,6 +10417,56 @@ 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;
+ }
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 741dee3a622..c42f20cb9fb 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -13314,8 +13314,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/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index b63886ae9dc..16e7502201b 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,10 @@
+2018-08-28 Julian Brown <julian@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-10-05 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
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/ChangeLog.omp b/libgomp/ChangeLog.omp
index f7d42c34b91..5ffb79b78b4 100644
--- a/libgomp/ChangeLog.omp
+++ b/libgomp/ChangeLog.omp
@@ -1,3 +1,15 @@
+2018-08-28 Julian Brown <julian@codesourcery.com>
+ Cesar Philippidis <cesar@codesourcery.com>
+
+ * 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.
+
2018-10-05 Nathan Sidwell <nathan@acm.org>
Tom de Vries <tdevries@suse.de>
Thomas Schwinge <thomas@codesourcery.com>
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
^ permalink raw reply [flat|nested] only message in thread
only message in thread, other threads:[~2021-05-13 16:09 UTC | newest]
Thread overview: (only message) (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2021-05-13 16:09 [gcc/devel/omp/gcc-11] Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828) Kwok Yeung
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).