This is the mail archive of the gcc-patches@gcc.gnu.org mailing list for the GCC project.


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]
Other format: [Raw text]

[PATCH, OpenACC] (1/2) Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828)


This patch implements support for array slices (with a non-zero base
element) declared on OpenACC data constructs. Any lexically-enclosed
parallel or kernels regions should "inherit" such mappings, e.g. if we
have:

#pragma acc data copy(arr[10:20])
{
#pragma acc parallel loop
  for (...) { ...arr[X]... }
}

the mapping for "arr" on the data construct takes precedence over the
default mapping behaviour for the parallel construct, which is to map
the whole array. (OpenACC 2.5, "2.5.1. Parallel Construct" and
elsewhere).

Tested with offloading to nvptx. (This patch differs in implementation
somewhat from the version on the gomp4, etc. branches.)

OK to apply?

Thanks,

Julian

2018-08-28  Julian Brown  <julian@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	PR middle-end/70828

	gcc/
	* gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash map.
	(new_omp_context): Initialise above.
	(delete_omp_context): Delete above.
	(gimplify_scan_omp_clauses): Scan for array mappings on data constructs,
	and record in above map.
	(gomp_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 omp 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-5.f90: New test.
>From 9123c4ddd701c40c3e85a0c6cd327066542b9e7a Mon Sep 17 00:00:00 2001
From: Julian Brown <julian@codesourcery.com>
Date: Thu, 16 Aug 2018 20:02:10 -0700
Subject: [PATCH 1/2] Inheritance of array sections on data constructs.

2018-08-28  Julian Brown  <julian@codesourcery.com>
	    Cesar Philippidis  <cesar@codesourcery.com>

	gcc/
	* gimplify.c (gimplify_omp_ctx): Add decl_data_clause hash map.
	(new_omp_context): Initialise above.
	(delete_omp_context): Delete above.
	(gimplify_scan_omp_clauses): Scan for array mappings on data constructs,
	and record in above map.
	(gomp_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 omp 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-5.f90: New test.
---
 gcc/gimplify.c                                     | 97 +++++++++++++++++++++-
 gcc/omp-low.c                                      |  7 +-
 gcc/testsuite/c-c++-common/goacc/acc-data-chain.c  | 24 ++++++
 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 ++++++
 .../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-5.f90   | 29 +++++++
 libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 | 24 ++++++
 11 files changed, 354 insertions(+), 5 deletions(-)
 create mode 100644 gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/pr70828.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/implicit_copy.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
 create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index dbd0f0e..d704aef 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -191,6 +191,7 @@ struct gimplify_omp_ctx
   bool target_map_scalars_firstprivate;
   bool target_map_pointers_as_0len_arrays;
   bool target_firstprivatize_array_bases;
+  hash_map<tree, std::pair<tree, tree> > *decl_data_clause;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -413,6 +414,7 @@ new_omp_context (enum omp_region_type region_type)
     c->default_kind = OMP_CLAUSE_DEFAULT_SHARED;
   else
     c->default_kind = OMP_CLAUSE_DEFAULT_UNSPECIFIED;
+  c->decl_data_clause = new hash_map<tree, std::pair<tree, tree> >;
 
   return c;
 }
@@ -425,6 +427,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);
 }
 
@@ -7793,8 +7796,21 @@ 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 nextc = OMP_CLAUSE_CHAIN (c);
+		if (nextc
+		    && OMP_CLAUSE_CODE (nextc) == OMP_CLAUSE_MAP
+		    && (OMP_CLAUSE_MAP_KIND (nextc)
+			  == GOMP_MAP_FIRSTPRIVATE_POINTER
+			|| OMP_CLAUSE_MAP_KIND (nextc) == GOMP_MAP_POINTER))
+		  {
+	            tree base_addr = OMP_CLAUSE_DECL (nextc);
+		    ctx->decl_data_clause->put (base_addr,
+		      std::make_pair (unshare_expr (c), unshare_expr (nextc)));
+		  }
+		if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
+		  break;
+	      }
 	      /* FALLTHRU */
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
@@ -8692,6 +8708,45 @@ 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 pair consisting of the mapping for
+   the data itself and for the pointer to the beginning of the data if present
+   in an outer data construct, or returns NULL otherwise.  */
+
+static std::pair<tree, tree> *
+gomp_needs_data_present (tree decl)
+{
+  gimplify_omp_ctx *ctx = NULL;
+
+  if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE
+      && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+      && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+	  || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE))
+    return NULL;
+
+  if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL
+      && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS)
+    return NULL;
+
+  decl = get_base_address (decl);
+
+  for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context)
+    {
+      std::pair<tree, tree> *ret;
+
+      if (ctx->region_type != ORT_ACC_DATA)
+	break;
+
+      if ((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.  */
 
@@ -8784,6 +8839,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;
+  std::pair<tree, tree> *mapping_ptr;
   if (private_debug)
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
@@ -8792,6 +8848,43 @@ 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)
+	   && (mapping_ptr = gomp_needs_data_present (decl)))
+    {
+      tree mapping = mapping_ptr->first;
+      tree pointer = mapping_ptr->second;
+
+      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));
+
+      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) = 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 fdabf67..be2bb73 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -8411,9 +8411,10 @@ 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))
-		      bias = lookup_decl (bias, ctx);
+		    tree bias = OMP_CLAUSE_SIZE (c), remapped_bias;
+		    if (DECL_P (bias)
+			&& (remapped_bias = maybe_lookup_decl (bias, ctx)))
+		      bias = remapped_bias;
 		    bias = fold_convert_loc (clause_loc, sizetype, bias);
 		    bias = fold_build1_loc (clause_loc, NEGATE_EXPR, sizetype,
 					    bias);
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 0000000..8a039be
--- /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 0000000..2e58120
--- /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 0000000..357114c
--- /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 0000000..4b6dbd7
--- /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 0000000..7a99f29
--- /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 0000000..22a9566
--- /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 0000000..ff17d10
--- /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-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
new file mode 100644
index 0000000..8a16e3d
--- /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.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
new file mode 100644
index 0000000..f87d232
--- /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
-- 
1.8.1.1


Index Nav: [Date Index] [Subject Index] [Author Index] [Thread Index]
Message Nav: [Date Prev] [Date Next] [Thread Prev] [Thread Next]