[08/13] OpenACC 2.6 deep copy: middle-end parts

Message ID 62e541dadce16450ba0da2bad66caf9bf443cd6e.1576648001.git.julian@codesourcery.com
State New
Headers show
Series
  • OpenACC 2.6 manual deep copy support
Related show

Commit Message

Julian Brown Dec. 18, 2019, 6:03 a.m.
This patch has been broken out of the "OpenACC 2.6 manual deep copy
support" patch, last posted here:

  https://gcc.gnu.org/ml/gcc-patches/2019-11/msg02376.html

This part contains the middle-end support for OpenACC 2.6 attach and
detach operations, either as standalone clauses or as "attach/detach"
actions triggered by other (data movement) clauses, as detailed in the
specification.

Tested alongside other patches in this series with offloading to
NVPTX. OK?

Thanks,

Julian

ChangeLog

	gcc/
	* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
	(insert_struct_comp_map): Support derived-type member mappings
	for arrays with descriptors which use GOMP_MAP_TO_PSET.  Support
	GOMP_MAP_ATTACH_DETACH.
	(gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA
	mappings.  Handle attach/detach clauses and component references.
	(gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
	attach/detach clauses.
	(gimplify_omp_target_update): Handle finalize for detach.
	* omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,
	GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.
	* tree-pretty-print.c (dump_omp_clause): Likewise, plus
	GOMP_MAP_ATTACH_DETACH.

	include/
	* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_ATTACH_DETACH.
---
 gcc/gimplify.c           | 232 ++++++++++++++++++++++++++++++++++-----
 gcc/omp-low.c            |   3 +
 gcc/tree-pretty-print.c  |  18 +++
 include/gomp-constants.h |   6 +-
 4 files changed, 229 insertions(+), 30 deletions(-)

-- 
2.23.0

Comments

Thomas Schwinge Dec. 21, 2019, 9:39 p.m. | #1
Hi!

On 2019-12-17T22:03:48-0800, Julian Brown <julian@codesourcery.com> wrote:
> This part contains the middle-end support for OpenACC 2.6 attach and

> detach operations, either as standalone clauses or as "attach/detach"

> actions triggered by other (data movement) clauses, as detailed in the

> specification.


As mentioned in <https://gcc.gnu.org/PR93026>, "that commit [r279626] is
doing more than just the OpenACC 2.6 manual deep copy changes; see the
'gcc/gimplify.c' changes related to the PR92929 discussion" etc. (see
<https://gcc.gnu.org/PR92929>).  See attached "[PR93026, PR92929] Adjust
'gfortran.dg/goacc/finalize-1.f' for r279626 changes"; committed to trunk
in r279700.


Grüße
 Thomas


> ChangeLog

>

> 	gcc/

> 	* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.

> 	(insert_struct_comp_map): Support derived-type member mappings

> 	for arrays with descriptors which use GOMP_MAP_TO_PSET.  Support

> 	GOMP_MAP_ATTACH_DETACH.

> 	(gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA

> 	mappings.  Handle attach/detach clauses and component references.

> 	(gimplify_adjust_omp_clauses_1): Skip adjustments for explicit

> 	attach/detach clauses.

> 	(gimplify_omp_target_update): Handle finalize for detach.

> 	* omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,

> 	GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.

> 	* tree-pretty-print.c (dump_omp_clause): Likewise, plus

> 	GOMP_MAP_ATTACH_DETACH.

>

> 	include/

> 	* gomp-constants.h (gomp_map_kind): Add GOMP_MAP_ATTACH_DETACH.

> ---

>  gcc/gimplify.c           | 232 ++++++++++++++++++++++++++++++++++-----

>  gcc/omp-low.c            |   3 +

>  gcc/tree-pretty-print.c  |  18 +++

>  include/gomp-constants.h |   6 +-

>  4 files changed, 229 insertions(+), 30 deletions(-)

>

> diff --git a/gcc/gimplify.c b/gcc/gimplify.c

> index e3088dcbe05..e3d5bc83c4f 100644

> --- a/gcc/gimplify.c

> +++ b/gcc/gimplify.c

> @@ -123,6 +123,10 @@ enum gimplify_omp_var_data

>    /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause.  */

>    GOVD_REDUCTION_INSCAN = 0x2000000,

>  

> +  /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for

> +     fields.  */

> +  GOVD_MAP_HAS_ATTACHMENTS = 8388608,

> +

>    GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE

>  			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR

>  			   | GOVD_LOCAL)

> @@ -8209,20 +8213,33 @@ insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,

>  			tree prev_node, tree *scp)

>  {

>    enum gomp_map_kind mkind

> -    = code == OMP_TARGET_EXIT_DATA ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;

> +    = (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)

> +      ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;

>  

>    tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);

>    tree cl = scp ? prev_node : c2;

>    OMP_CLAUSE_SET_MAP_KIND (c2, mkind);

>    OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));

>    OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;

> -  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);

> +  if (OMP_CLAUSE_CHAIN (prev_node) != c

> +      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP

> +      && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))

> +	  == GOMP_MAP_TO_PSET))

> +    OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));

> +  else

> +    OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);

>    if (struct_node)

>      OMP_CLAUSE_CHAIN (struct_node) = c2;

>  

>    /* We might need to create an additional mapping if we have a reference to a

> -     pointer (in C++).  */

> -  if (OMP_CLAUSE_CHAIN (prev_node) != c)

> +     pointer (in C++).  Don't do this if we have something other than a

> +     GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET.  */

> +  if (OMP_CLAUSE_CHAIN (prev_node) != c

> +      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP

> +      && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))

> +	   == GOMP_MAP_ALWAYS_POINTER)

> +	  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))

> +	      == GOMP_MAP_ATTACH_DETACH)))

>      {

>        tree c4 = OMP_CLAUSE_CHAIN (prev_node);

>        tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);

> @@ -8329,6 +8346,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>    struct gimplify_omp_ctx *ctx, *outer_ctx;

>    tree c;

>    hash_map<tree, tree> *struct_map_to_clause = NULL;

> +  hash_set<tree> *struct_deref_set = NULL;

>    tree *prev_list_p = NULL, *orig_list_p = list_p;

>    int handled_depend_iterators = -1;

>    int nowait = -1;

> @@ -8731,8 +8749,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  	    case OMP_TARGET_DATA:

>  	    case OMP_TARGET_ENTER_DATA:

>  	    case OMP_TARGET_EXIT_DATA:

> -	    case OACC_ENTER_DATA:

> -	    case OACC_EXIT_DATA:

>  	    case OACC_HOST_DATA:

>  	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER

>  		  || (OMP_CLAUSE_MAP_KIND (c)

> @@ -8741,6 +8757,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  		   mapped, but not the pointer to it.  */

>  		remove = true;

>  	      break;

> +	    case OACC_ENTER_DATA:

> +	    case OACC_EXIT_DATA:

> +	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER

> +		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET

> +		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER

> +		  || (OMP_CLAUSE_MAP_KIND (c)

> +		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))

> +		remove = true;

> +	      break;

>  	    default:

>  	      break;

>  	    }

> @@ -8814,7 +8839,35 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  		  pd = &TREE_OPERAND (decl, 0);

>  		  decl = TREE_OPERAND (decl, 0);

>  		}

> -	      if (TREE_CODE (decl) == COMPONENT_REF)

> +	      bool indir_p = false;

> +	      tree orig_decl = decl;

> +	      tree decl_ref = NULL_TREE;

> +	      if ((region_type & ORT_ACC) != 0

> +		  && TREE_CODE (*pd) == COMPONENT_REF

> +		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH

> +		  && code != OACC_UPDATE)

> +		{

> +		  while (TREE_CODE (decl) == COMPONENT_REF)

> +		    {

> +		      decl = TREE_OPERAND (decl, 0);

> +		      if ((TREE_CODE (decl) == MEM_REF

> +			   && integer_zerop (TREE_OPERAND (decl, 1)))

> +			  || INDIRECT_REF_P (decl))

> +			{

> +			  indir_p = true;

> +			  decl = TREE_OPERAND (decl, 0);

> +			}

> +		      if (TREE_CODE (decl) == INDIRECT_REF

> +			  && DECL_P (TREE_OPERAND (decl, 0))

> +			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))

> +			      == REFERENCE_TYPE))

> +			{

> +			  decl_ref = decl;

> +			  decl = TREE_OPERAND (decl, 0);

> +			}

> +		    }

> +		}

> +	      else if (TREE_CODE (decl) == COMPONENT_REF)

>  		{

>  		  while (TREE_CODE (decl) == COMPONENT_REF)

>  		    decl = TREE_OPERAND (decl, 0);

> @@ -8824,13 +8877,76 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  			  == REFERENCE_TYPE))

>  		    decl = TREE_OPERAND (decl, 0);

>  		}

> +	      if (decl != orig_decl && DECL_P (decl) && indir_p)

> +		{

> +		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH

> +							     : GOMP_MAP_ATTACH;

> +		  /* We have a dereference of a struct member.  Make this an

> +		     attach/detach operation, and ensure the base pointer is

> +		     mapped as a FIRSTPRIVATE_POINTER.  */

> +		  OMP_CLAUSE_SET_MAP_KIND (c, k);

> +		  flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;

> +		  tree next_clause = OMP_CLAUSE_CHAIN (c);

> +		  if (k == GOMP_MAP_ATTACH

> +		      && code != OACC_ENTER_DATA

> +		      && (!next_clause

> +			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)

> +			   || (OMP_CLAUSE_MAP_KIND (next_clause)

> +			       != GOMP_MAP_POINTER)

> +			   || OMP_CLAUSE_DECL (next_clause) != decl)

> +		      && (!struct_deref_set

> +			  || !struct_deref_set->contains (decl)))

> +		    {

> +		      if (!struct_deref_set)

> +			struct_deref_set = new hash_set<tree> ();

> +		      /* As well as the attach, we also need a

> +			 FIRSTPRIVATE_POINTER clause to properly map the

> +			 pointer to the struct base.  */

> +		      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),

> +						  OMP_CLAUSE_MAP);

> +		      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);

> +		      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)

> +			= 1;

> +		      tree charptr_zero

> +			= build_int_cst (build_pointer_type (char_type_node),

> +					 0);

> +		      OMP_CLAUSE_DECL (c2)

> +			= build2 (MEM_REF, char_type_node,

> +				  decl_ref ? decl_ref : decl, charptr_zero);

> +		      OMP_CLAUSE_SIZE (c2) = size_zero_node;

> +		      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),

> +						  OMP_CLAUSE_MAP);

> +		      OMP_CLAUSE_SET_MAP_KIND (c3,

> +					       GOMP_MAP_FIRSTPRIVATE_POINTER);

> +		      OMP_CLAUSE_DECL (c3) = decl;

> +		      OMP_CLAUSE_SIZE (c3) = size_zero_node;

> +		      tree mapgrp = *prev_list_p;

> +		      *prev_list_p = c2;

> +		      OMP_CLAUSE_CHAIN (c3) = mapgrp;

> +		      OMP_CLAUSE_CHAIN (c2) = c3;

> +

> +		      struct_deref_set->add (decl);

> +		    }

> +		  goto do_add_decl;

> +		}

> +	      /* An "attach/detach" operation on an update directive should

> +		 behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that

> +		 unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER

> +		 depends on the previous mapping.  */

> +	      if (code == OACC_UPDATE

> +		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)

> +		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);

>  	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)

>  		  == GS_ERROR)

>  		{

>  		  remove = true;

>  		  break;

>  		}

> -	      if (DECL_P (decl))

> +	      if (DECL_P (decl)

> +		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET

> +		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH

> +		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH

> +		  && code != OACC_UPDATE)

>  		{

>  		  if (error_operand_p (decl))

>  		    {

> @@ -8851,7 +8967,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  		      break;

>  		    }

>  

> -		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)

> +		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER

> +		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)

>  		    {

>  		      /* Error recovery.  */

>  		      if (prev_list_p == NULL)

> @@ -8884,20 +9001,47 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);

>  		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)

>  			      == GOMP_MAP_ALWAYS_POINTER);

> +		  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)

> +					== GOMP_MAP_ATTACH_DETACH);

> +		  bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH

> +				|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH;

> +		  bool has_attachments = false;

> +		  /* For OpenACC, pointers in structs should trigger an

> +		     attach action.  */

> +		  if (attach_detach && (region_type & ORT_ACC) != 0)

> +		    {

> +		      /* Turn a GOMP_MAP_ATTACH_DETACH clause into a

> +			 GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we

> +			 have detected a case that needs a GOMP_MAP_STRUCT

> +			 mapping added.  */

> +		      gomp_map_kind k

> +			= (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH

> +						   : GOMP_MAP_ATTACH;

> +		      OMP_CLAUSE_SET_MAP_KIND (c, k);

> +		      has_attachments = true;

> +		    }

>  		  if (n == NULL || (n->value & GOVD_MAP) == 0)

>  		    {

>  		      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),

>  						 OMP_CLAUSE_MAP);

> -		      OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);

> +		      gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT

> +					       : GOMP_MAP_STRUCT;

> +

> +		      OMP_CLAUSE_SET_MAP_KIND (l, k);

>  		      if (base_ref)

>  			OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);

>  		      else

>  			OMP_CLAUSE_DECL (l) = decl;

> -		      OMP_CLAUSE_SIZE (l) = size_int (1);

> +		      OMP_CLAUSE_SIZE (l)

> +			= (!attach

> +			   ? size_int (1)

> +			   : DECL_P (OMP_CLAUSE_DECL (l))

> +			   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))

> +			   : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));

>  		      if (struct_map_to_clause == NULL)

>  			struct_map_to_clause = new hash_map<tree, tree>;

>  		      struct_map_to_clause->put (decl, l);

> -		      if (ptr)

> +		      if (ptr || attach_detach)

>  			{

>  			  insert_struct_comp_map (code, c, l, *prev_list_p,

>  						  NULL);

> @@ -8923,23 +9067,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  			  OMP_CLAUSE_CHAIN (l) = c2;

>  			}

>  		      flags = GOVD_MAP | GOVD_EXPLICIT;

> -		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)

> +		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))

> +			  || ptr

> +			  || attach_detach)

>  			flags |= GOVD_SEEN;

> +		      if (has_attachments)

> +			flags |= GOVD_MAP_HAS_ATTACHMENTS;

>  		      goto do_add_decl;

>  		    }

> -		  else

> +		  else if (struct_map_to_clause)

>  		    {

>  		      tree *osc = struct_map_to_clause->get (decl);

>  		      tree *sc = NULL, *scp = NULL;

> -		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)

> +		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))

> +			  || ptr

> +			  || attach_detach)

>  			n->value |= GOVD_SEEN;

>  		      sc = &OMP_CLAUSE_CHAIN (*osc);

>  		      if (*sc != c

>  			  && (OMP_CLAUSE_MAP_KIND (*sc)

> -			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) 

> +			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))

>  			sc = &OMP_CLAUSE_CHAIN (*sc);

> +		      /* Here "prev_list_p" is the end of the inserted

> +			 alloc/release nodes after the struct node, OSC.  */

>  		      for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))

> -			if (ptr && sc == prev_list_p)

> +			if ((ptr || attach_detach) && sc == prev_list_p)

>  			  break;

>  			else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))

>  				 != COMPONENT_REF

> @@ -8992,7 +9144,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  				|| (known_eq (offset1, offsetn)

>  				    && maybe_lt (bitpos1, bitposn)))

>  			      {

> -				if (ptr)

> +				if (ptr || attach_detach)

>  				  scp = sc;

>  				else

>  				  break;

> @@ -9000,10 +9152,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  			  }

>  		      if (remove)

>  			break;

> -		      OMP_CLAUSE_SIZE (*osc)

> -			= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),

> -				      size_one_node);

> -		      if (ptr)

> +		      if (!attach)

> +			OMP_CLAUSE_SIZE (*osc)

> +			  = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),

> +					size_one_node);

> +		      if (ptr || attach_detach)

>  			{

>  			  tree cl = insert_struct_comp_map (code, c, NULL,

>  							    *prev_list_p, scp);

> @@ -9033,11 +9186,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>  		}

>  	      if (!remove

>  		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER

> +		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH

> +		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET

>  		  && OMP_CLAUSE_CHAIN (c)

>  		  && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP

> -		  && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))

> -		      == GOMP_MAP_ALWAYS_POINTER))

> +		  && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))

> +		       == GOMP_MAP_ALWAYS_POINTER)

> +		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))

> +			  == GOMP_MAP_ATTACH_DETACH)

> +		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))

> +			  == GOMP_MAP_TO_PSET)))

>  		prev_list_p = list_p;

> +

>  	      break;

>  	    }

>  	  flags = GOVD_MAP | GOVD_EXPLICIT;

> @@ -9561,6 +9721,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,

>    gimplify_omp_ctxp = ctx;

>    if (struct_map_to_clause)

>      delete struct_map_to_clause;

> +  if (struct_deref_set)

> +    delete struct_deref_set;

>  }

>  

>  /* Return true if DECL is a candidate for shared to firstprivate

> @@ -9708,6 +9870,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)

>      return 0;

>    if ((flags & GOVD_SEEN) == 0)

>      return 0;

> +  if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)

> +    return 0;

>    if (flags & GOVD_DEBUG_PRIVATE)

>      {

>        gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);

> @@ -12762,8 +12926,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)

>  	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),

>  			       OMP_CLAUSE_FINALIZE))

>      {

> -      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"

> -	 semantics apply to all mappings of this OpenACC directive.  */

> +      /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and

> +	 GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply

> +	 to all mappings of this OpenACC directive.  */

>        bool finalize_marked = false;

>        for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))

>  	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)

> @@ -12777,10 +12942,19 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)

>  	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);

>  	      finalize_marked = true;

>  	      break;

> +	    case GOMP_MAP_DETACH:

> +	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);

> +	      finalize_marked = true;

> +	      break;

> +	    case GOMP_MAP_STRUCT:

> +	    case GOMP_MAP_FORCE_PRESENT:

> +	      /* Skip over an initial struct or force_present mapping.  */

> +	      break;

>  	    default:

> -	      /* Check consistency: libgomp relies on the very first data

> -		 mapping clause being marked, so make sure we did that before

> -		 any other mapping clauses.  */

> +	      /* Check consistency: libgomp relies on the very first

> +		 non-struct, non-force-present data mapping clause being

> +		 marked, so make sure we did that before any other mapping

> +		 clauses.  */

>  	      gcc_assert (finalize_marked);

>  	      break;

>  	    }

> diff --git a/gcc/omp-low.c b/gcc/omp-low.c

> index d422c205836..3eb7815449a 100644

> --- a/gcc/omp-low.c

> +++ b/gcc/omp-low.c

> @@ -11439,6 +11439,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)

>  	  case GOMP_MAP_FORCE_DEVICEPTR:

>  	  case GOMP_MAP_DEVICE_RESIDENT:

>  	  case GOMP_MAP_LINK:

> +	  case GOMP_MAP_ATTACH:

> +	  case GOMP_MAP_DETACH:

> +	  case GOMP_MAP_FORCE_DETACH:

>  	    gcc_assert (is_gimple_omp_oacc (stmt));

>  	    break;

>  	  default:

> diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c

> index 1cf7a912133..379858d0f1f 100644

> --- a/gcc/tree-pretty-print.c

> +++ b/gcc/tree-pretty-print.c

> @@ -849,6 +849,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)

>  	case GOMP_MAP_LINK:

>  	  pp_string (pp, "link");

>  	  break;

> +	case GOMP_MAP_ATTACH:

> +	  pp_string (pp, "attach");

> +	  break;

> +	case GOMP_MAP_DETACH:

> +	  pp_string (pp, "detach");

> +	  break;

> +	case GOMP_MAP_FORCE_DETACH:

> +	  pp_string (pp, "force_detach");

> +	  break;

> +	case GOMP_MAP_ATTACH_DETACH:

> +	  pp_string (pp, "attach_detach");

> +	  break;

>  	default:

>  	  gcc_unreachable ();

>  	}

> @@ -870,6 +882,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)

>  	    case GOMP_MAP_TO_PSET:

>  	      pp_string (pp, " [pointer set, len: ");

>  	      break;

> +	    case GOMP_MAP_ATTACH:

> +	    case GOMP_MAP_DETACH:

> +	    case GOMP_MAP_FORCE_DETACH:

> +	    case GOMP_MAP_ATTACH_DETACH:

> +	      pp_string (pp, " [bias: ");

> +	      break;

>  	    default:

>  	      pp_string (pp, " [len: ");

>  	      break;

> diff --git a/include/gomp-constants.h b/include/gomp-constants.h

> index e8bd52e81bd..f40d6069582 100644

> --- a/include/gomp-constants.h

> +++ b/include/gomp-constants.h

> @@ -142,7 +142,11 @@ enum gomp_map_kind

>      /* Do not map, but pointer assign a pointer instead.  */

>      GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1),

>      /* Do not map, but pointer assign a reference instead.  */

> -    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2)

> +    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2),

> +    /* An attach or detach operation.  Rewritten to the appropriate type during

> +       gimplification, depending on directive (i.e. "enter data" or

> +       parallel/kernels region vs. "exit data").  */

> +    GOMP_MAP_ATTACH_DETACH =		(GOMP_MAP_LAST | 3)

>    };

>  

>  #define GOMP_MAP_COPY_TO_P(X) \
From 0c1f5b1c22e0c0c3dd0b93697de2235af7e4adfa Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>

Date: Sat, 21 Dec 2019 21:32:36 +0000
Subject: [PATCH] [PR93026, PR92929] Adjust 'gfortran.dg/goacc/finalize-1.f'
 for r279626 changes

	gcc/testsuite/
	PR fortran/93026
	PR middle-end/92929
	* gfortran.dg/goacc/finalize-1.f: Adjust.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279700 138bc75d-0d04-0410-961f-82ee72b054a4
---
 gcc/testsuite/ChangeLog                      | 6 ++++++
 gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++--
 2 files changed, 8 insertions(+), 2 deletions(-)

diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index bbd9131e5cd..219ff3b9284 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-12-21  Thomas Schwinge  <thomas@codesourcery.com>
+
+	PR fortran/93026
+	PR middle-end/92929
+	* gfortran.dg/goacc/finalize-1.f: Adjust.
+
 2019-12-21  Harald Anlauf  <anlauf@gmx.de>
 
 	PR fortran/91661
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index ca642156e9f..1e2e3e94b8a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -21,7 +21,7 @@
 
 !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_r)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -33,5 +33,5 @@
 
 !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
-- 
2.17.1

Patch

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index e3088dcbe05..e3d5bc83c4f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -123,6 +123,10 @@  enum gimplify_omp_var_data
   /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause.  */
   GOVD_REDUCTION_INSCAN = 0x2000000,
 
+  /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
+     fields.  */
+  GOVD_MAP_HAS_ATTACHMENTS = 8388608,
+
   GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
 			   | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
 			   | GOVD_LOCAL)
@@ -8209,20 +8213,33 @@  insert_struct_comp_map (enum tree_code code, tree c, tree struct_node,
 			tree prev_node, tree *scp)
 {
   enum gomp_map_kind mkind
-    = code == OMP_TARGET_EXIT_DATA ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+    = (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+      ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
 
   tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
   tree cl = scp ? prev_node : c2;
   OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
   OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));
   OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;
-  OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
+  if (OMP_CLAUSE_CHAIN (prev_node) != c
+      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
+      && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	  == GOMP_MAP_TO_PSET))
+    OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));
+  else
+    OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
   if (struct_node)
     OMP_CLAUSE_CHAIN (struct_node) = c2;
 
   /* We might need to create an additional mapping if we have a reference to a
-     pointer (in C++).  */
-  if (OMP_CLAUSE_CHAIN (prev_node) != c)
+     pointer (in C++).  Don't do this if we have something other than a
+     GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET.  */
+  if (OMP_CLAUSE_CHAIN (prev_node) != c
+      && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
+      && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	   == GOMP_MAP_ALWAYS_POINTER)
+	  || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+	      == GOMP_MAP_ATTACH_DETACH)))
     {
       tree c4 = OMP_CLAUSE_CHAIN (prev_node);
       tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -8329,6 +8346,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   struct gimplify_omp_ctx *ctx, *outer_ctx;
   tree c;
   hash_map<tree, tree> *struct_map_to_clause = NULL;
+  hash_set<tree> *struct_deref_set = NULL;
   tree *prev_list_p = NULL, *orig_list_p = list_p;
   int handled_depend_iterators = -1;
   int nowait = -1;
@@ -8731,8 +8749,6 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
-	    case OACC_ENTER_DATA:
-	    case OACC_EXIT_DATA:
 	    case OACC_HOST_DATA:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
@@ -8741,6 +8757,15 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		   mapped, but not the pointer to it.  */
 		remove = true;
 	      break;
+	    case OACC_ENTER_DATA:
+	    case OACC_EXIT_DATA:
+	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		  || (OMP_CLAUSE_MAP_KIND (c)
+		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		remove = true;
+	      break;
 	    default:
 	      break;
 	    }
@@ -8814,7 +8839,35 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
-	      if (TREE_CODE (decl) == COMPONENT_REF)
+	      bool indir_p = false;
+	      tree orig_decl = decl;
+	      tree decl_ref = NULL_TREE;
+	      if ((region_type & ORT_ACC) != 0
+		  && TREE_CODE (*pd) == COMPONENT_REF
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+		  && code != OACC_UPDATE)
+		{
+		  while (TREE_CODE (decl) == COMPONENT_REF)
+		    {
+		      decl = TREE_OPERAND (decl, 0);
+		      if ((TREE_CODE (decl) == MEM_REF
+			   && integer_zerop (TREE_OPERAND (decl, 1)))
+			  || INDIRECT_REF_P (decl))
+			{
+			  indir_p = true;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		      if (TREE_CODE (decl) == INDIRECT_REF
+			  && DECL_P (TREE_OPERAND (decl, 0))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == REFERENCE_TYPE))
+			{
+			  decl_ref = decl;
+			  decl = TREE_OPERAND (decl, 0);
+			}
+		    }
+		}
+	      else if (TREE_CODE (decl) == COMPONENT_REF)
 		{
 		  while (TREE_CODE (decl) == COMPONENT_REF)
 		    decl = TREE_OPERAND (decl, 0);
@@ -8824,13 +8877,76 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  == REFERENCE_TYPE))
 		    decl = TREE_OPERAND (decl, 0);
 		}
+	      if (decl != orig_decl && DECL_P (decl) && indir_p)
+		{
+		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+							     : GOMP_MAP_ATTACH;
+		  /* We have a dereference of a struct member.  Make this an
+		     attach/detach operation, and ensure the base pointer is
+		     mapped as a FIRSTPRIVATE_POINTER.  */
+		  OMP_CLAUSE_SET_MAP_KIND (c, k);
+		  flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
+		  tree next_clause = OMP_CLAUSE_CHAIN (c);
+		  if (k == GOMP_MAP_ATTACH
+		      && code != OACC_ENTER_DATA
+		      && (!next_clause
+			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
+			   || (OMP_CLAUSE_MAP_KIND (next_clause)
+			       != GOMP_MAP_POINTER)
+			   || OMP_CLAUSE_DECL (next_clause) != decl)
+		      && (!struct_deref_set
+			  || !struct_deref_set->contains (decl)))
+		    {
+		      if (!struct_deref_set)
+			struct_deref_set = new hash_set<tree> ();
+		      /* As well as the attach, we also need a
+			 FIRSTPRIVATE_POINTER clause to properly map the
+			 pointer to the struct base.  */
+		      tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+		      OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
+			= 1;
+		      tree charptr_zero
+			= build_int_cst (build_pointer_type (char_type_node),
+					 0);
+		      OMP_CLAUSE_DECL (c2)
+			= build2 (MEM_REF, char_type_node,
+				  decl_ref ? decl_ref : decl, charptr_zero);
+		      OMP_CLAUSE_SIZE (c2) = size_zero_node;
+		      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+						  OMP_CLAUSE_MAP);
+		      OMP_CLAUSE_SET_MAP_KIND (c3,
+					       GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      OMP_CLAUSE_DECL (c3) = decl;
+		      OMP_CLAUSE_SIZE (c3) = size_zero_node;
+		      tree mapgrp = *prev_list_p;
+		      *prev_list_p = c2;
+		      OMP_CLAUSE_CHAIN (c3) = mapgrp;
+		      OMP_CLAUSE_CHAIN (c2) = c3;
+
+		      struct_deref_set->add (decl);
+		    }
+		  goto do_add_decl;
+		}
+	      /* An "attach/detach" operation on an update directive should
+		 behave as a GOMP_MAP_ALWAYS_POINTER.  Beware that
+		 unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
+		 depends on the previous mapping.  */
+	      if (code == OACC_UPDATE
+		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+		OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
 	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
 		  == GS_ERROR)
 		{
 		  remove = true;
 		  break;
 		}
-	      if (DECL_P (decl))
+	      if (DECL_P (decl)
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+		  && code != OACC_UPDATE)
 		{
 		  if (error_operand_p (decl))
 		    {
@@ -8851,7 +8967,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      break;
 		    }
 
-		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+		  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+		      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		    {
 		      /* Error recovery.  */
 		      if (prev_list_p == NULL)
@@ -8884,20 +9001,47 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		    = splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
 		  bool ptr = (OMP_CLAUSE_MAP_KIND (c)
 			      == GOMP_MAP_ALWAYS_POINTER);
+		  bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
+					== GOMP_MAP_ATTACH_DETACH);
+		  bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+				|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH;
+		  bool has_attachments = false;
+		  /* For OpenACC, pointers in structs should trigger an
+		     attach action.  */
+		  if (attach_detach && (region_type & ORT_ACC) != 0)
+		    {
+		      /* Turn a GOMP_MAP_ATTACH_DETACH clause into a
+			 GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we
+			 have detected a case that needs a GOMP_MAP_STRUCT
+			 mapping added.  */
+		      gomp_map_kind k
+			= (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+						   : GOMP_MAP_ATTACH;
+		      OMP_CLAUSE_SET_MAP_KIND (c, k);
+		      has_attachments = true;
+		    }
 		  if (n == NULL || (n->value & GOVD_MAP) == 0)
 		    {
 		      tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 						 OMP_CLAUSE_MAP);
-		      OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+		      gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT
+					       : GOMP_MAP_STRUCT;
+
+		      OMP_CLAUSE_SET_MAP_KIND (l, k);
 		      if (base_ref)
 			OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
 		      else
 			OMP_CLAUSE_DECL (l) = decl;
-		      OMP_CLAUSE_SIZE (l) = size_int (1);
+		      OMP_CLAUSE_SIZE (l)
+			= (!attach
+			   ? size_int (1)
+			   : DECL_P (OMP_CLAUSE_DECL (l))
+			   ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+			   : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
 		      if (struct_map_to_clause == NULL)
 			struct_map_to_clause = new hash_map<tree, tree>;
 		      struct_map_to_clause->put (decl, l);
-		      if (ptr)
+		      if (ptr || attach_detach)
 			{
 			  insert_struct_comp_map (code, c, l, *prev_list_p,
 						  NULL);
@@ -8923,23 +9067,31 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  OMP_CLAUSE_CHAIN (l) = c2;
 			}
 		      flags = GOVD_MAP | GOVD_EXPLICIT;
-		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+			  || ptr
+			  || attach_detach)
 			flags |= GOVD_SEEN;
+		      if (has_attachments)
+			flags |= GOVD_MAP_HAS_ATTACHMENTS;
 		      goto do_add_decl;
 		    }
-		  else
+		  else if (struct_map_to_clause)
 		    {
 		      tree *osc = struct_map_to_clause->get (decl);
 		      tree *sc = NULL, *scp = NULL;
-		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+		      if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+			  || ptr
+			  || attach_detach)
 			n->value |= GOVD_SEEN;
 		      sc = &OMP_CLAUSE_CHAIN (*osc);
 		      if (*sc != c
 			  && (OMP_CLAUSE_MAP_KIND (*sc)
-			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) 
+			      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 			sc = &OMP_CLAUSE_CHAIN (*sc);
+		      /* Here "prev_list_p" is the end of the inserted
+			 alloc/release nodes after the struct node, OSC.  */
 		      for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
-			if (ptr && sc == prev_list_p)
+			if ((ptr || attach_detach) && sc == prev_list_p)
 			  break;
 			else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
 				 != COMPONENT_REF
@@ -8992,7 +9144,7 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 				|| (known_eq (offset1, offsetn)
 				    && maybe_lt (bitpos1, bitposn)))
 			      {
-				if (ptr)
+				if (ptr || attach_detach)
 				  scp = sc;
 				else
 				  break;
@@ -9000,10 +9152,11 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  }
 		      if (remove)
 			break;
-		      OMP_CLAUSE_SIZE (*osc)
-			= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
-				      size_one_node);
-		      if (ptr)
+		      if (!attach)
+			OMP_CLAUSE_SIZE (*osc)
+			  = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+					size_one_node);
+		      if (ptr || attach_detach)
 			{
 			  tree cl = insert_struct_comp_map (code, c, NULL,
 							    *prev_list_p, scp);
@@ -9033,11 +9186,18 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      if (!remove
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
+		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
 		  && OMP_CLAUSE_CHAIN (c)
 		  && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
-		  && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
-		      == GOMP_MAP_ALWAYS_POINTER))
+		  && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+		       == GOMP_MAP_ALWAYS_POINTER)
+		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			  == GOMP_MAP_ATTACH_DETACH)
+		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+			  == GOMP_MAP_TO_PSET)))
 		prev_list_p = list_p;
+
 	      break;
 	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -9561,6 +9721,8 @@  gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
   gimplify_omp_ctxp = ctx;
   if (struct_map_to_clause)
     delete struct_map_to_clause;
+  if (struct_deref_set)
+    delete struct_deref_set;
 }
 
 /* Return true if DECL is a candidate for shared to firstprivate
@@ -9708,6 +9870,8 @@  gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
     return 0;
   if ((flags & GOVD_SEEN) == 0)
     return 0;
+  if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
+    return 0;
   if (flags & GOVD_DEBUG_PRIVATE)
     {
       gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
@@ -12762,8 +12926,9 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	   && omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
 			       OMP_CLAUSE_FINALIZE))
     {
-      /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
-	 semantics apply to all mappings of this OpenACC directive.  */
+      /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
+	 GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
+	 to all mappings of this OpenACC directive.  */
       bool finalize_marked = false;
       for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
 	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -12777,10 +12942,19 @@  gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
 	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
 	      finalize_marked = true;
 	      break;
+	    case GOMP_MAP_DETACH:
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
+	      finalize_marked = true;
+	      break;
+	    case GOMP_MAP_STRUCT:
+	    case GOMP_MAP_FORCE_PRESENT:
+	      /* Skip over an initial struct or force_present mapping.  */
+	      break;
 	    default:
-	      /* Check consistency: libgomp relies on the very first data
-		 mapping clause being marked, so make sure we did that before
-		 any other mapping clauses.  */
+	      /* Check consistency: libgomp relies on the very first
+		 non-struct, non-force-present data mapping clause being
+		 marked, so make sure we did that before any other mapping
+		 clauses.  */
 	      gcc_assert (finalize_marked);
 	      break;
 	    }
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index d422c205836..3eb7815449a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11439,6 +11439,9 @@  lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FORCE_DEVICEPTR:
 	  case GOMP_MAP_DEVICE_RESIDENT:
 	  case GOMP_MAP_LINK:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
+	  case GOMP_MAP_FORCE_DETACH:
 	    gcc_assert (is_gimple_omp_oacc (stmt));
 	    break;
 	  default:
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 1cf7a912133..379858d0f1f 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -849,6 +849,18 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	case GOMP_MAP_LINK:
 	  pp_string (pp, "link");
 	  break;
+	case GOMP_MAP_ATTACH:
+	  pp_string (pp, "attach");
+	  break;
+	case GOMP_MAP_DETACH:
+	  pp_string (pp, "detach");
+	  break;
+	case GOMP_MAP_FORCE_DETACH:
+	  pp_string (pp, "force_detach");
+	  break;
+	case GOMP_MAP_ATTACH_DETACH:
+	  pp_string (pp, "attach_detach");
+	  break;
 	default:
 	  gcc_unreachable ();
 	}
@@ -870,6 +882,12 @@  dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
 	    case GOMP_MAP_TO_PSET:
 	      pp_string (pp, " [pointer set, len: ");
 	      break;
+	    case GOMP_MAP_ATTACH:
+	    case GOMP_MAP_DETACH:
+	    case GOMP_MAP_FORCE_DETACH:
+	    case GOMP_MAP_ATTACH_DETACH:
+	      pp_string (pp, " [bias: ");
+	      break;
 	    default:
 	      pp_string (pp, " [len: ");
 	      break;
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index e8bd52e81bd..f40d6069582 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -142,7 +142,11 @@  enum gomp_map_kind
     /* Do not map, but pointer assign a pointer instead.  */
     GOMP_MAP_FIRSTPRIVATE_POINTER =	(GOMP_MAP_LAST | 1),
     /* Do not map, but pointer assign a reference instead.  */
-    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2)
+    GOMP_MAP_FIRSTPRIVATE_REFERENCE =	(GOMP_MAP_LAST | 2),
+    /* An attach or detach operation.  Rewritten to the appropriate type during
+       gimplification, depending on directive (i.e. "enter data" or
+       parallel/kernels region vs. "exit data").  */
+    GOMP_MAP_ATTACH_DETACH =		(GOMP_MAP_LAST | 3)
   };
 
 #define GOMP_MAP_COPY_TO_P(X) \