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]

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


On Tue, 4 Dec 2018 15:02:15 +0100
Jakub Jelinek <jakub@redhat.com> wrote:

> On Tue, Aug 28, 2018 at 03:19:19PM -0400, Julian Brown wrote:
> > 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.  
> 
> > --- 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>
> > >;  
> 
> Not really happy about creating this unconditionally.  Can you leave
> it NULL by default and only initialize for contexts where it will be
> needed?
> 
> > @@ -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)));  
> 
> Don't like the wrapping here, can you just split it up:
> 		    std::pair<tree, tree> p
> 		      = std::make_pair (unshare_expr (c),
> 					unshare_expr (nextc));
> 		    ctx->decl_data_clause->put (base_addr, p);
> or similar?
> 
> > +
> > +static std::pair<tree, tree> *
> > +gomp_needs_data_present (tree decl)  
> 
> Would be helpful to have acc/oacc in the function name.
> > +{
> > +  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;  
> 
> And move this test to the top.
> 
> > --- 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);  
> 
> This is shared with OpenMP and must be conditionalized for OpenACC
> only.

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.

Re-tested with offloading to nvptx. OK?

Julian

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-5.f90: New test.              
        * testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.

Reviewed-by: Jakub Jelinek  <jakub@redhat.com>
commit 390adf97cfdde951ed1e82fc54d77e34130c70b8
Author: Julian Brown <julian@codesourcery.com>
Date:   Thu Aug 16 20:02:10 2018 -0700

    Inheritance of array sections on data constructs.
    
    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-5.f90: New test.
    	* testsuite/libgomp.oacc-fortran/pr70828-6.f90: New test.

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 509fc2f..b6a9bfc 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -176,6 +176,17 @@ struct gimplify_ctx
   unsigned in_switch_expr : 1;
 };
 
+/* 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;
@@ -191,6 +202,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, oacc_array_mapping_info> *decl_data_clause;
 };
 
 static struct gimplify_ctx *gimplify_ctxp;
@@ -413,6 +425,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 = NULL;
 
   return c;
 }
@@ -425,6 +438,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);
 }
 
@@ -7795,8 +7810,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:
@@ -8695,6 +8743,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.  */
 
@@ -8787,6 +8875,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))
@@ -8795,6 +8884,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 fdabf67..84c9a88 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -8411,8 +8411,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 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-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
new file mode 100644
index 0000000..e99c364
--- /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 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

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