diff mbox series

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

Message ID 20180828151919.576c636c@squid.athome
State New
Headers show
Series [OpenACC] (1/2) Fix implicit mapping for array slices on lexically-enclosing data constructs (PR70828) | expand

Commit Message

Julian Brown Aug. 28, 2018, 7:19 p.m. UTC
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.

Comments

Jakub Jelinek Dec. 4, 2018, 2:02 p.m. UTC | #1
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.

	Jakub
Julian Brown Dec. 5, 2018, 9:10 p.m. UTC | #2
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
Julian Brown June 9, 2019, 5:35 p.m. UTC | #3
On Wed, 5 Dec 2018 21:10:45 +0000
Julian Brown <julian@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  <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/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 mbox series

Patch

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