[PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 1458 bytes --]

These are the middle-end gimplifier and omp-low changes.
Compiler testcases are also included in this patch.

For attach/detach clauses, I'm currently using the clause tree expression
itself as the key for lookup, to solve the "same-decl" problem when
multiple clauses have the same OMP_CLAUSE_DECL. This is just a special
case right now, have yet to see if this can expand to more general use
between all map clauses.

Thanks,
Chung-Lin

	gcc/
	* gimplify.c (is_or_contains_p): New static helper function.
	(omp_target_reorder_clauses): New function.
	(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
	reorder clause list according to OpenMP 5.0 rules. Add handling of
	GOMP_MAP_ATTACH_DETACH for OpenMP cases.
	* omp-low.c (is_omp_target): New static helper function.
	(scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH
	for OpenMP cases.
	(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
	OpenMP cases.

	gcc/testsuite/
	* c-c++-common/goacc/finalize-1.c: Adjust gimple scanning.
         * c-c++-common/goacc/mdc-1.c: Likewise.
         * c-c++-common/goacc/struct-enter-exit-data-1.c: Likewise.
         * gfortran.dg/goacc/attach-descriptor.f90: Likewise.
         * gfortran.dg/goacc/finalize-1.f: Likewise.
         * c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
         * gfortran.dg/gomp/map-2.f90: Likewise.
         * c-c++-common/gomp/map-5.c: New testcase.

[-- Attachment #2: omp5-tgtmapping.02.midend-testsuite.patch --]
[-- Type: text/plain, Size: 41502 bytes --]

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 23d0e2511f7..0ad141c5b3f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8350,14 +8350,126 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   /* Set *BASE_REF if BASE was a dereferenced reference variable.  */
   if (base_ref && orig_base != base)
     *base_ref = orig_base;
 
   return base;
 }
 
+/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
+
+static bool
+is_or_contains_p (tree expr, tree base_ptr)
+{
+  while (expr != base_ptr)
+    if (TREE_CODE (base_ptr) == COMPONENT_REF)
+      base_ptr = TREE_OPERAND (base_ptr, 0);
+    else
+      break;
+  return expr == base_ptr;
+}
+
+/* Implement OpenMP 5.x map ordering rules for target directives. There are
+   several rules, and with some level of ambiguity, hopefully we can at least
+   collect the complexity here in one place.  */
+
+static void
+omp_target_reorder_clauses (tree *list_p)
+{
+  vec<tree> clauses = vNULL;
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    clauses.safe_push (*cp);
+
+  /* Collect refs to alloc/release/delete maps.  */
+  vec<tree> ard = vNULL;
+  for (unsigned int i = 0; i < clauses.length (); i++)
+    if (OMP_CLAUSE_CODE (clauses[i]) == OMP_CLAUSE_MAP
+	&& (OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_ALLOC
+	    || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_RELEASE
+	    || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_DELETE))
+      {
+	ard.safe_push (clauses[i]);
+	clauses[i] = NULL_TREE;
+
+	unsigned int j;
+	for (j = i + 1; j < clauses.length (); j++)
+	  {
+	    /* Any associated pointer type maps should move along.  */
+	    tree nc = clauses[j];
+	    if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+		    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH
+		    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
+		    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
+		    || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+	      {
+		ard.safe_push (nc);
+		clauses[j] = NULL_TREE;
+	      }
+	    else
+	      break;
+	  }
+	i = j - 1;
+      }
+
+  tree *cp = list_p;
+  for (unsigned int i = 0; i < clauses.length (); i++)
+    if (clauses[i])
+      {
+	*cp = clauses[i];
+	cp = &OMP_CLAUSE_CHAIN (clauses[i]);
+      }
+  for (unsigned int i = 0; i < ard.length (); i++)
+    {
+      *cp = ard[i];
+      cp = &OMP_CLAUSE_CHAIN (ard[i]);
+    }
+  *cp = NULL_TREE;
+
+  /* OpenMP 5.0 requires that pointer variables are mapped before
+     its use as a base-pointer.  */
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+      {
+	tree decl = OMP_CLAUSE_DECL (*cp);
+	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+	if ((k == GOMP_MAP_ALLOC
+	     || k == GOMP_MAP_TO
+	     || k == GOMP_MAP_FROM
+	     || k == GOMP_MAP_TOFROM)
+	    && (TREE_CODE (decl) == INDIRECT_REF
+		|| TREE_CODE (decl) == MEM_REF))
+	  {
+	    tree base_ptr = TREE_OPERAND (decl, 0);
+	    STRIP_TYPE_NOPS (base_ptr);
+	    for (tree *cp2 = &OMP_CLAUSE_CHAIN (*cp); *cp2;
+		 cp2 = &OMP_CLAUSE_CHAIN (*cp2))
+	      if (OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP)
+		{
+		  tree decl2 = OMP_CLAUSE_DECL (*cp2);
+		  gomp_map_kind k2 = OMP_CLAUSE_MAP_KIND (*cp2);
+		  if ((k2 == GOMP_MAP_ALLOC
+		       || k2 == GOMP_MAP_TO
+		       || k2 == GOMP_MAP_FROM
+		       || k2 == GOMP_MAP_TOFROM)
+		      && is_or_contains_p (decl2, base_ptr))
+		    {
+		      /* Move *cp2 to before *cp.  */
+		      tree c = *cp2;
+		      *cp2 = OMP_CLAUSE_CHAIN (c);
+		      OMP_CLAUSE_CHAIN (c) = *cp;
+		      *cp = c;
+		      if (*cp2 == NULL_TREE)
+			break;
+		    }
+		}
+	  }
+      }
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
 static void
 gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			   enum omp_region_type region_type,
 			   enum tree_code code)
@@ -8391,14 +8503,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
       case OACC_PARALLEL:
       case OACC_KERNELS:
 	ctx->target_firstprivatize_array_bases = true;
       default:
 	break;
       }
 
+  if (code == OMP_TARGET
+      || code == OMP_TARGET_DATA
+      || code == OMP_TARGET_ENTER_DATA
+      || code == OMP_TARGET_EXIT_DATA)
+    omp_target_reorder_clauses (list_p);
+
   while ((c = *list_p) != NULL)
     {
       bool remove = false;
       bool notice_outer = true;
       const char *check_non_private = NULL;
       unsigned int flags;
       tree decl;
@@ -8831,23 +8949,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			     NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
 	    {
 	      remove = true;
 	      break;
 	    }
 	  else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		    || (OMP_CLAUSE_MAP_KIND (c)
-			== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
 	    {
 	      OMP_CLAUSE_SIZE (c)
 		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
 					   false);
-	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
-				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	      if ((region_type & ORT_TARGET) != 0)
+		omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
+
 	  if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
 	      if (TREE_CODE (d) == ARRAY_REF)
 		{
 		  while (TREE_CODE (d) == ARRAY_REF)
 		    d = TREE_OPERAND (d, 0);
@@ -8864,25 +8985,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		{
 		  pd = &TREE_OPERAND (decl, 0);
 		  decl = TREE_OPERAND (decl, 0);
 		}
 	      bool indir_p = false;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
-	      if ((region_type & ORT_ACC) != 0
+	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 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))
+		      if (((TREE_CODE (decl) == MEM_REF
+			    && integer_zerop (TREE_OPERAND (decl, 1)))
+			   || INDIRECT_REF_P (decl))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == POINTER_TYPE))
 			{
 			  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)))
@@ -8901,24 +9024,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		      && DECL_P (TREE_OPERAND (decl, 0))
 		      && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
 			  == 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;
+		  gomp_map_kind k
+		    = ((code == OACC_EXIT_DATA || code == OMP_TARGET_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
+		      && code != OMP_TARGET_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)))
@@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      /* 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)
 		  && 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)
+		  && code != OACC_UPDATE
+		  && code != OMP_TARGET_UPDATE)
 		{
 		  if (error_operand_p (decl))
 		    {
 		      remove = true;
 		      break;
 		    }
 
@@ -9030,23 +9150,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  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)
+		  if (attach_detach
+		      && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
+			  || code == OMP_TARGET_ENTER_DATA
+			  || code == OMP_TARGET_EXIT_DATA))
+
 		    {
 		      /* 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;
+			= ((code == OACC_EXIT_DATA || code == OMP_TARGET_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);
@@ -9134,41 +9258,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			    tree base
 			      = extract_base_bit_offset (sc_decl, NULL,
 							 &bitposn, &offsetn);
 			    if (base != decl)
 			      break;
 			    if (scp)
 			      continue;
-			    tree d1 = OMP_CLAUSE_DECL (*sc);
-			    tree d2 = OMP_CLAUSE_DECL (c);
-			    while (TREE_CODE (d1) == ARRAY_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    while (TREE_CODE (d2) == ARRAY_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    if (TREE_CODE (d1) == INDIRECT_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    if (TREE_CODE (d2) == INDIRECT_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    while (TREE_CODE (d1) == COMPONENT_REF)
-			      if (TREE_CODE (d2) == COMPONENT_REF
-				  && TREE_OPERAND (d1, 1)
-				     == TREE_OPERAND (d2, 1))
-				{
+			    if (! (code == OMP_TARGET
+				   || code == OMP_TARGET_DATA
+				   || code == OMP_TARGET_ENTER_DATA
+				   || code == OMP_TARGET_EXIT_DATA))
+			      {
+				/* This duplicate checking code is currently only
+				   enabled for OpenACC.  */
+				tree d1 = OMP_CLAUSE_DECL (*sc);
+				tree d2 = OMP_CLAUSE_DECL (c);
+				while (TREE_CODE (d1) == ARRAY_REF)
 				  d1 = TREE_OPERAND (d1, 0);
+				while (TREE_CODE (d2) == ARRAY_REF)
 				  d2 = TREE_OPERAND (d2, 0);
-				}
-			      else
-				break;
-			    if (d1 == d2)
-			      {
-				error_at (OMP_CLAUSE_LOCATION (c),
-					  "%qE appears more than once in map "
-					  "clauses", OMP_CLAUSE_DECL (c));
-				remove = true;
-				break;
+				if (TREE_CODE (d1) == INDIRECT_REF)
+				  d1 = TREE_OPERAND (d1, 0);
+				if (TREE_CODE (d2) == INDIRECT_REF)
+				  d2 = TREE_OPERAND (d2, 0);
+				while (TREE_CODE (d1) == COMPONENT_REF)
+				  if (TREE_CODE (d2) == COMPONENT_REF
+				      && TREE_OPERAND (d1, 1)
+				      == TREE_OPERAND (d2, 1))
+				    {
+				      d1 = TREE_OPERAND (d1, 0);
+				      d2 = TREE_OPERAND (d2, 0);
+				    }
+				  else
+				    break;
+				if (d1 == d2)
+				  {
+				    error_at (OMP_CLAUSE_LOCATION (c),
+					      "%qE appears more than once in map "
+					      "clauses", OMP_CLAUSE_DECL (c));
+				    remove = true;
+				    break;
+				  }
 			      }
 			    if (maybe_lt (offset1, offsetn)
 				|| (known_eq (offset1, offsetn)
 				    && maybe_lt (bitpos1, bitposn)))
 			      {
 				if (ptr || attach_detach)
 				  scp = sc;
@@ -9222,18 +9354,68 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			  == GOMP_MAP_ATTACH_DETACH)
 		      || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
 			  == GOMP_MAP_TO_PSET)))
 		prev_list_p = list_p;
 
 	      break;
 	    }
+	  else
+	    {
+	      /* DECL_P (decl) == true  */
+	      tree *sc;
+	      if (struct_map_to_clause
+		  && (sc = struct_map_to_clause->get (decl)) != NULL
+		  && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
+		  && decl == OMP_CLAUSE_DECL (*sc))
+		{
+		  /* We have found a map of the whole structure after a
+		     leading GOMP_MAP_STRUCT has been created, so refill the
+		     leading clause into a map of the whole structure
+		     variable, and remove the current one.
+		     TODO: we should be able to remove some maps of the
+		     following structure element maps if they are of
+		     compatible TO/FROM/ALLOC type.  */
+		  OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
+		  OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
+		  remove = true;
+		  break;
+		}
+	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
+
+	  if ((code == OMP_TARGET
+	       || code == OMP_TARGET_DATA
+	       || code == OMP_TARGET_ENTER_DATA
+	       || code == OMP_TARGET_EXIT_DATA)
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	    {
+	      for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
+		   octx = octx->outer_context)
+		{
+		  splay_tree_node n
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) OMP_CLAUSE_DECL (c));
+		  /* If this is contained in an outer OpenMP region as a
+		     firstprivate value, remove the attach/detach.  */
+		  if (n && (n->value & GOVD_FIRSTPRIVATE))
+		    {
+		      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      goto do_add;
+		    }
+		}
+
+	      enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
+					     ? GOMP_MAP_DETACH
+					     : GOMP_MAP_ATTACH);
+	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+	    }
+
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
 	  if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
 	    {
 	      tree deps = OMP_CLAUSE_DECL (c);
 	      while (deps && TREE_CODE (deps) == TREE_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 53efe5f750c..8d50774384a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -210,14 +210,29 @@ is_oacc_kernels (omp_context *ctx)
 {
   enum gimple_code outer_type = gimple_code (ctx->stmt);
   return ((outer_type == GIMPLE_OMP_TARGET)
 	  && (gimple_omp_target_kind (ctx->stmt)
 	      == GF_OMP_TARGET_KIND_OACC_KERNELS));
 }
 
+/* Return true if STMT corresponds to an OpenMP target region.  */
+static bool
+is_omp_target (gimple *stmt)
+{
+  if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+    {
+      int kind = gimple_omp_target_kind (stmt);
+      return (kind == GF_OMP_TARGET_KIND_REGION
+	      || kind == GF_OMP_TARGET_KIND_DATA
+	      || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+	      || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+    }
+  return false;
+}
+
 /* If DECL is the artificial dummy VAR_DECL created for non-static
    data member privatization, return the underlying "this" parameter,
    otherwise return NULL.  */
 
 tree
 omp_member_access_dummy_var (tree decl)
 {
@@ -1342,15 +1357,17 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	     don't need to be copied, the receiver side will use them
 	     directly.  However, global variables with "omp declare target link"
 	     attribute need to be copied.  Or when ALWAYS modifier is used.  */
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && DECL_P (decl)
 	      && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
 		   && (OMP_CLAUSE_MAP_KIND (c)
-		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
 		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable
 	      && !lookup_attribute ("omp declare target link",
@@ -1362,14 +1379,48 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      /* Ignore GOMP_MAP_POINTER kind for arrays in regions that are
 		 not offloaded; there is nothing to map for those.  */
 	      if (!is_gimple_omp_offloaded (ctx->stmt)
 		  && !POINTER_TYPE_P (TREE_TYPE (decl))
 		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && DECL_P (decl)
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt))
+	    {
+	      /* If this is an offloaded region, an attach operation should
+		 only exist when the pointer variable is mapped in a prior
+		 clause.  */
+	      if (is_gimple_omp_offloaded (ctx->stmt))
+		gcc_assert
+		  (maybe_lookup_decl (decl, ctx)
+		   || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+		       && lookup_attribute ("omp declare target",
+					    DECL_ATTRIBUTES (decl))));
+
+	      /* By itself, attach/detach is generated as part of pointer
+		 variable mapping and should not create new variables in the
+		 offloaded region, however sender refs for it must be created
+		 for its address to be passed to the runtime.  */
+	      tree field
+		= build_decl (OMP_CLAUSE_LOCATION (c),
+			      FIELD_DECL, NULL_TREE, ptr_type_node);
+	      SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+	      insert_field_into_struct (ctx->record_type, field);
+	      /* To not clash with a map of the pointer variable itself,
+		 attach/detach maps have their field looked up by the *clause*
+		 tree expression, not the decl.  */
+	      gcc_assert (!splay_tree_lookup (ctx->field_map,
+					      (splay_tree_key) c));
+	      splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+				 (splay_tree_value) field);
+	      break;
+	    }
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
 		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
 	    {
 	      if (TREE_CODE (decl) == COMPONENT_REF
 		  || (TREE_CODE (decl) == INDIRECT_REF
@@ -1601,14 +1652,19 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
 		   && (OMP_CLAUSE_MAP_KIND (c)
 		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
 		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable)
 	    break;
+	  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt)
+	      && !is_gimple_omp_offloaded (ctx->stmt))
+	    break;
 	  if (DECL_P (decl))
 	    {
 	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		   || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 		  && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
 		  && !COMPLETE_TYPE_P (TREE_TYPE (decl)))
 		{
@@ -11405,26 +11461,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_ALWAYS_TO:
 	  case GOMP_MAP_ALWAYS_FROM:
 	  case GOMP_MAP_ALWAYS_TOFROM:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_STRUCT:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
 	    break;
 	  case GOMP_MAP_IF_PRESENT:
 	  case GOMP_MAP_FORCE_ALLOC:
 	  case GOMP_MAP_FORCE_TO:
 	  case GOMP_MAP_FORCE_FROM:
 	  case GOMP_MAP_FORCE_TOFROM:
 	  case GOMP_MAP_FORCE_PRESENT:
 	  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:
 	    gcc_unreachable ();
 	  }
 #endif
@@ -11471,14 +11527,24 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		x = build_simple_mem_ref (x);
 		SET_DECL_VALUE_EXPR (new_var, x);
 		DECL_HAS_VALUE_EXPR_P (new_var) = 1;
 	      }
 	    continue;
 	  }
 
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	    && is_omp_target (stmt))
+	  {
+	    gcc_assert (maybe_lookup_field (c, ctx));
+	    map_cnt++;
+	    continue;
+	  }
+
 	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
 	/* Don't remap compute constructs' reduction variables, because the
 	   intermediate result must be local to each gang.  */
 	if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 			   && OMP_CLAUSE_MAP_IN_REDUCTION (c)))
@@ -11703,22 +11769,36 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		  {
 		    tree ovar2 = DECL_VALUE_EXPR (ovar);
 		    gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
 		    ovar2 = TREE_OPERAND (ovar2, 0);
 		    gcc_assert (DECL_P (ovar2));
 		    ovar = ovar2;
 		  }
-		if (!maybe_lookup_field (ovar, ctx))
+		if (!maybe_lookup_field (ovar, ctx)
+		    && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			 && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+			     || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
 		  continue;
 	      }
 
 	    talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
 	    if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
 	      talign = DECL_ALIGN_UNIT (ovar);
-	    if (nc)
+
+	    if (nc
+		&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		&& is_omp_target (stmt))
+	      {
+		var = lookup_decl_in_outer_ctx (ovar, ctx);
+		x = build_sender_ref (c, ctx);
+		gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+	      }
+	    else if (nc)
 	      {
 		var = lookup_decl_in_outer_ctx (ovar, ctx);
 		x = build_sender_ref (ovar, ctx);
 
 		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
 		    && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 3d64b2e7cb3..679b0505e19 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -17,21 +17,21 @@ void f ()
 
 #pragma acc exit data finalize delete (del_f)
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } } */
 
 #pragma acc exit data finalize delete (del_f_p[2:5])
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
-   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
 
 #pragma acc exit data copyout (cpo_r)
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
 
 #pragma acc exit data copyout (cpo_f) finalize
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:cpo_f\\);$" 1 "original" } }
    { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } } */
 
 #pragma acc exit data copyout (cpo_f_p[4:10]) finalize
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
-   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\)$" 1 "gimple" } } */
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 337c1f7cc77..839269eb62b 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -40,15 +40,15 @@ t1 ()
 
 #pragma acc exit data detach(a) finalize
 #pragma acc exit data detach(s.a) finalize
   }
 }
 
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:\*.*z.? .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:\*.*s\.a.? .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
index df405e448b2..9f702ba76f2 100644
--- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -16,12 +16,12 @@ struct str {
 
 void
 test (int *b, int *c, int *e)
 {
   struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
 
 #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
-  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*.*s\.b.? \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*.*s\.c.? \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
 
 #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
-  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*.*s\.b.? \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*.*s\.c.? \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index ded1d74ccde..bbc8fb4e32b 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -9,46 +9,46 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
     bar (&q);
   #pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more than once in data clauses" } */
     bar (p);
   #pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more than once in data clauses" } */
     bar (p);
   #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
     bar (p);
-  #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
+  #pragma omp target map (p) , map (p[0])
     bar (p);
   #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
     bar (&q);
   #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
     bar (p);
-  #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map (t.r)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t.r)
     bar (&t.r);
   #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
   #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
-  #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t)
     bar (t.s);
-  #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map(t.s[0])
     bar (t.s);
   #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
-  #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t.s[2])
     bar (t.s);
-  #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[0:2]) map (t.t[4:6])
     bar (t.t);
-  #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[i:j]) map (t.t[k:l])
     bar (t.t);
   #pragma omp target map (t.s[0]) map (t.r)
     bar (t.s);
   #pragma omp target map (t.r) ,map (t.s[0])
     bar (t.s);
   #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
-    bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
+    bar (t.s);
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c
new file mode 100644
index 00000000000..1d9d9252864
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-5.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+  /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses.  */
+  int a, b, c;
+  #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
+  #pragma omp target exit data map(from:a) map(release:b) map(from:c)
+
+  #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
+  a = b = c = 1;
+
+  #pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
+  #pragma omp target exit data map(from:a) map(delete:b) map(from:c)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 373bdcb2114..c5ac06943eb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -8,22 +8,22 @@ program att
   end type t
   type(t) :: myvar
   integer, target :: tarr(10)
   integer, pointer :: myptr(:)
 
   !$acc enter data attach(myvar%arr2, myptr)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   !$acc exit data detach(myvar%arr2, myptr)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   ! Test valid usage and processing of the finalize clause.
   !$acc exit data detach(myvar%arr2, myptr) finalize
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } }
   ! For array-descriptor detaches, we no longer generate a "release" mapping
   ! for the pointed-to data for gimplify.c to turn into "delete".  Make sure
   ! the mapping still isn't there.
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
 
 end program att
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index a7788580819..0ff2e471180 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -17,21 +17,21 @@
 
 !$ACC EXIT DATA FINALIZE DELETE (del_f)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:del_f\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) 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" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
 
 !$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_f\\) finalize;$" 1 "original" } }
 ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
 
 !$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: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
index 73c4f5a87d0..79bab726dea 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
@@ -1,6 +1,6 @@
 type t
   integer :: i
 end type t
 type(t) v
-!$omp target enter data map(to:v%i, v%i)  ! { dg-error "appears more than once in map clauses" }
+!$omp target enter data map(to:v%i, v%i)
 end

Re: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Jakub Jelinek]

On Tue, Sep 01, 2020 at 09:16:48PM +0800, Chung-Lin Tang wrote:
> --- a/gcc/gimplify.c
> +++ b/gcc/gimplify.c
> @@ -8350,14 +8350,126 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
>    /* Set *BASE_REF if BASE was a dereferenced reference variable.  */
>    if (base_ref && orig_base != base)
>      *base_ref = orig_base;
>  
>    return base;
>  }
>  
> +/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
> +
> +static bool
> +is_or_contains_p (tree expr, tree base_ptr)
> +{
> +  while (expr != base_ptr)
> +    if (TREE_CODE (base_ptr) == COMPONENT_REF)
> +      base_ptr = TREE_OPERAND (base_ptr, 0);
> +    else
> +      break;
> +  return expr == base_ptr;
> +}
> +
> +/* Implement OpenMP 5.x map ordering rules for target directives. There are
> +   several rules, and with some level of ambiguity, hopefully we can at least
> +   collect the complexity here in one place.  */
> +
> +static void
> +omp_target_reorder_clauses (tree *list_p)
> +{

So, first of all, are you convinced we can sort just the explicit clauses
and leave out the (later on) implicitly added ones?
If it is possible, sure, it will be easier (because we later on need to deal
with the GOMP_MAP_STRUCT sorting too).

> +  vec<tree> clauses = vNULL;

Isn't this a memory leak?  Nothing frees the vector.  Perhaps better
  auto_vec<tree, 32> clauses;

> +  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
> +    clauses.safe_push (*cp);

The rest of the function deals only with OMP_CLAUSE_MAP clauses, wouldn't it
be better to just push to the vec those clauses and keep other clauses just
in *list_p chain?

> +  /* Collect refs to alloc/release/delete maps.  */
> +  vec<tree> ard = vNULL;

Again, auto_vec<tree, 32> ard;

> +  tree *cp = list_p;
> +  for (unsigned int i = 0; i < clauses.length (); i++)
> +    if (clauses[i])
> +      {
> +	*cp = clauses[i];
> +	cp = &OMP_CLAUSE_CHAIN (clauses[i]);
> +      }
> +  for (unsigned int i = 0; i < ard.length (); i++)
> +    {
> +      *cp = ard[i];
> +      cp = &OMP_CLAUSE_CHAIN (ard[i]);
> +    }
> +  *cp = NULL_TREE;
> +
> +  /* OpenMP 5.0 requires that pointer variables are mapped before
> +     its use as a base-pointer.  */
> +  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
> +    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
> +      {
> +	tree decl = OMP_CLAUSE_DECL (*cp);
> +	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
> +	if ((k == GOMP_MAP_ALLOC
> +	     || k == GOMP_MAP_TO
> +	     || k == GOMP_MAP_FROM
> +	     || k == GOMP_MAP_TOFROM)

What about the *ALWAYS* kinds?

> +	    && (TREE_CODE (decl) == INDIRECT_REF
> +		|| TREE_CODE (decl) == MEM_REF))
> +	  {
> +	    tree base_ptr = TREE_OPERAND (decl, 0);
> +	    STRIP_TYPE_NOPS (base_ptr);
> +	    for (tree *cp2 = &OMP_CLAUSE_CHAIN (*cp); *cp2;
> +		 cp2 = &OMP_CLAUSE_CHAIN (*cp2))
> +	      if (OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP)
> +		{
> +		  tree decl2 = OMP_CLAUSE_DECL (*cp2);
> +		  gomp_map_kind k2 = OMP_CLAUSE_MAP_KIND (*cp2);
> +		  if ((k2 == GOMP_MAP_ALLOC
> +		       || k2 == GOMP_MAP_TO
> +		       || k2 == GOMP_MAP_FROM
> +		       || k2 == GOMP_MAP_TOFROM)

Again.

This is O(n^2) too, but due to the is_or_contains_p I'm not sure
if we can avoid it.  Perhaps sort the clauses by uid of the base expressions
and deal with those separately.  Maybe let's ignore it for now.

> @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>  	      /* 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;
> -		}

So what gimplifies those now?

> +			    if (! (code == OMP_TARGET
> +				   || code == OMP_TARGET_DATA
> +				   || code == OMP_TARGET_ENTER_DATA
> +				   || code == OMP_TARGET_EXIT_DATA))
> +			      {

Isn't this just if ((region_type & ORT_ACC) == 0) ?  Or do we want
it for target update too?  Though, we wouldn't talk about more than once in
map clauses then because target update doesn't have those.


	Jakub

Re: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 5004 bytes --]

On 2020/10/13 9:31 PM, Jakub Jelinek wrote:
>> +/* Implement OpenMP 5.x map ordering rules for target directives. There are
>> +   several rules, and with some level of ambiguity, hopefully we can at least
>> +   collect the complexity here in one place.  */
>> +
>> +static void
>> +omp_target_reorder_clauses (tree *list_p)
>> +{
> So, first of all, are you convinced we can sort just the explicit clauses
> and leave out the (later on) implicitly added ones?
> If it is possible, sure, it will be easier (because we later on need to deal
> with the GOMP_MAP_STRUCT sorting too).

Yeah, there will probably be more cases to handle later, and possibly sinking
the call to omp_target_reorder_clauses till after the main handling in
gimplify_scan_omp_clauses. But the current routine handles a straightforward
set of cases, which can be grown upon later.

>> +  vec<tree> clauses = vNULL;
> Isn't this a memory leak?  Nothing frees the vector.  Perhaps better
>    auto_vec<tree, 32> clauses;
> 
>> +  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
>> +    clauses.safe_push (*cp);
> The rest of the function deals only with OMP_CLAUSE_MAP clauses, wouldn't it
> be better to just push to the vec those clauses and keep other clauses just
> in *list_p chain?
> 
>> +  /* Collect refs to alloc/release/delete maps.  */
>> +  vec<tree> ard = vNULL;
> Again, auto_vec<tree, 32> ard;

Thanks for catching this. I'm now using auto_vec now.

>> +  tree *cp = list_p;
>> +  for (unsigned int i = 0; i < clauses.length (); i++)
>> +    if (clauses[i])
>> +      {
>> +	*cp = clauses[i];
>> +	cp = &OMP_CLAUSE_CHAIN (clauses[i]);
>> +      }
>> +  for (unsigned int i = 0; i < ard.length (); i++)
>> +    {
>> +      *cp = ard[i];
>> +      cp = &OMP_CLAUSE_CHAIN (ard[i]);
>> +    }
>> +  *cp = NULL_TREE;
>> +
>> +  /* OpenMP 5.0 requires that pointer variables are mapped before
>> +     its use as a base-pointer.  */
>> +  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
>> +    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
>> +      {
>> +	tree decl = OMP_CLAUSE_DECL (*cp);
>> +	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
>> +	if ((k == GOMP_MAP_ALLOC
>> +	     || k == GOMP_MAP_TO
>> +	     || k == GOMP_MAP_FROM
>> +	     || k == GOMP_MAP_TOFROM)
> What about the*ALWAYS*  kinds?

Adjustment done, plus re-written so only one pass of this checking is done.

>> +	    && (TREE_CODE (decl) == INDIRECT_REF
>> +		|| TREE_CODE (decl) == MEM_REF))
>> +	  {
>> +	    tree base_ptr = TREE_OPERAND (decl, 0);
>> +	    STRIP_TYPE_NOPS (base_ptr);
>> +	    for (tree *cp2 = &OMP_CLAUSE_CHAIN (*cp); *cp2;
>> +		 cp2 = &OMP_CLAUSE_CHAIN (*cp2))
>> +	      if (OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP)
>> +		{
>> +		  tree decl2 = OMP_CLAUSE_DECL (*cp2);
>> +		  gomp_map_kind k2 = OMP_CLAUSE_MAP_KIND (*cp2);
>> +		  if ((k2 == GOMP_MAP_ALLOC
>> +		       || k2 == GOMP_MAP_TO
>> +		       || k2 == GOMP_MAP_FROM
>> +		       || k2 == GOMP_MAP_TOFROM)
> Again.
> 
> This is O(n^2) too, but due to the is_or_contains_p I'm not sure
> if we can avoid it.  Perhaps sort the clauses by uid of the base expressions
> and deal with those separately.  Maybe let's ignore it for now.

I re-wrote most of omp_target_reorder_clauses to be more efficient. The O(n^2)
issues should be fixed now.

>> @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>>   	      /* 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;
>> -		}
> So what gimplifies those now?

They're gimplified somewhere during omp-low now.
(some gimplify scan testcases were adjusted to accommodate this change)

I don't remember the exact case I encountered, but there were some issues with gimplified
expressions inside the map clauses making some later checking more difficult. I haven't seen
any negative effect of this modification so far.

>> +			    if (! (code == OMP_TARGET
>> +				   || code == OMP_TARGET_DATA
>> +				   || code == OMP_TARGET_ENTER_DATA
>> +				   || code == OMP_TARGET_EXIT_DATA))
>> +			      {
> Isn't this just if ((region_type & ORT_ACC) == 0) ?  Or do we want
> it for target update too?  Though, we wouldn't talk about more than once in
> map clauses then because target update doesn't have those.

It's actually "(region_type & ORT_ACC) != 0", which I now use in the patch.
I originally intended to be careful and only pick those four OpenMP target constructs
to conditionalize on, but so far using the above test works without regressions.

Updated patch attached.

Thanks,
Chung-Lin

[-- Attachment #2: omp5-tgtmapping.02.midend-testsuite.v2.patch --]
[-- Type: text/plain, Size: 33800 bytes --]

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 29f385c9368..d5048d140d8 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8364,6 +8364,113 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   return base;
 }
 
+/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
+
+static bool
+is_or_contains_p (tree expr, tree base_ptr)
+{
+  while (expr != base_ptr)
+    if (TREE_CODE (base_ptr) == COMPONENT_REF)
+      base_ptr = TREE_OPERAND (base_ptr, 0);
+    else
+      break;
+  return expr == base_ptr;
+}
+
+/* Implement OpenMP 5.x map ordering rules for target directives. There are
+   several rules, and with some level of ambiguity, hopefully we can at least
+   collect the complexity here in one place.  */
+
+static void
+omp_target_reorder_clauses (tree *list_p)
+{
+  /* Collect refs to alloc/release/delete maps.  */
+  auto_vec<tree, 32> ard;
+  tree *cp = list_p;
+  while (*cp != NULL_TREE)
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+	&& (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
+	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
+	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
+      {
+	/* Unlink cp and push to ard.  */
+	tree c = *cp;
+	tree nc = OMP_CLAUSE_CHAIN (c);
+	*cp = nc;
+	ard.safe_push (c);
+
+	/* Any associated pointer type maps should also move along.  */
+	while (*cp != NULL_TREE
+	       && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+	       && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
+	  {
+	    c = *cp;
+	    nc = OMP_CLAUSE_CHAIN (c);
+	    *cp = nc;
+	    ard.safe_push (c);
+	  }
+      }
+    else
+      cp = &OMP_CLAUSE_CHAIN (*cp);
+
+  /* Link alloc/release/delete maps to the end of list.  */
+  for (unsigned int i = 0; i < ard.length (); i++)
+    {
+      *cp = ard[i];
+      cp = &OMP_CLAUSE_CHAIN (ard[i]);
+    }
+  *cp = NULL_TREE;
+
+  /* OpenMP 5.0 requires that pointer variables are mapped before
+     its use as a base-pointer.  */
+  auto_vec<tree *, 32> atf;
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+      {
+	/* Collect alloc, to, from, to/from clause tree pointers.  */
+	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+	if (k == GOMP_MAP_ALLOC
+	    || k == GOMP_MAP_TO
+	    || k == GOMP_MAP_FROM
+	    || k == GOMP_MAP_TOFROM
+	    || k == GOMP_MAP_ALWAYS_TO
+	    || k == GOMP_MAP_ALWAYS_FROM
+	    || k == GOMP_MAP_ALWAYS_TOFROM)
+	  atf.safe_push (cp);
+      }
+
+  for (unsigned int i = 0; i < atf.length (); i++)
+    if (atf[i])
+      {
+	tree *cp = atf[i];
+	tree decl = OMP_CLAUSE_DECL (*cp);
+	if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
+	  {
+	    tree base_ptr = TREE_OPERAND (decl, 0);
+	    STRIP_TYPE_NOPS (base_ptr);
+	    for (unsigned int j = i + 1; j < atf.length (); j++)
+	      {
+		tree *cp2 = atf[j];
+		tree decl2 = OMP_CLAUSE_DECL (*cp2);
+		if (is_or_contains_p (decl2, base_ptr))
+		  {
+		    /* Move *cp2 to before *cp.  */
+		    tree c = *cp2;
+		    *cp2 = OMP_CLAUSE_CHAIN (c);
+		    OMP_CLAUSE_CHAIN (c) = *cp;
+		    *cp = c;
+		    atf[j] = NULL;
+		  }
+	      }
+	  }
+      }
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
@@ -8405,6 +8512,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	break;
       }
 
+  if (code == OMP_TARGET
+      || code == OMP_TARGET_DATA
+      || code == OMP_TARGET_ENTER_DATA
+      || code == OMP_TARGET_EXIT_DATA)
+    omp_target_reorder_clauses (list_p);
+
   while ((c = *list_p) != NULL)
     {
       bool remove = false;
@@ -8845,15 +8958,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		    || (OMP_CLAUSE_MAP_KIND (c)
-			== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
 	    {
 	      OMP_CLAUSE_SIZE (c)
 		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
 					   false);
-	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
-				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	      if ((region_type & ORT_TARGET) != 0)
+		omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
+
 	  if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
@@ -8878,7 +8994,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      bool indir_p = false;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
-	      if ((region_type & ORT_ACC) != 0
+	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
 		  && TREE_CODE (*pd) == COMPONENT_REF
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
 		  && code != OACC_UPDATE)
@@ -8886,9 +9002,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  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))
+		      if (((TREE_CODE (decl) == MEM_REF
+			    && integer_zerop (TREE_OPERAND (decl, 1)))
+			   || INDIRECT_REF_P (decl))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == POINTER_TYPE))
 			{
 			  indir_p = true;
 			  decl = TREE_OPERAND (decl, 0);
@@ -8915,8 +9033,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      if (decl != orig_decl && DECL_P (decl) && indir_p)
 		{
-		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
-							     : GOMP_MAP_ATTACH;
+		  gomp_map_kind k
+		    = ((code == OACC_EXIT_DATA || code == OMP_TARGET_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.  */
@@ -8925,6 +9044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  tree next_clause = OMP_CLAUSE_CHAIN (c);
 		  if (k == GOMP_MAP_ATTACH
 		      && code != OACC_ENTER_DATA
+		      && code != OMP_TARGET_ENTER_DATA
 		      && (!next_clause
 			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
 			   || (OMP_CLAUSE_MAP_KIND (next_clause)
@@ -8972,17 +9092,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      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)
 		  && 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)
+		  && code != OACC_UPDATE
+		  && code != OMP_TARGET_UPDATE)
 		{
 		  if (error_operand_p (decl))
 		    {
@@ -9044,15 +9159,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  bool has_attachments = false;
 		  /* For OpenACC, pointers in structs should trigger an
 		     attach action.  */
-		  if (attach_detach && (region_type & ORT_ACC) != 0)
+		  if (attach_detach
+		      && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
+			  || code == OMP_TARGET_ENTER_DATA
+			  || code == OMP_TARGET_EXIT_DATA))
+
 		    {
 		      /* 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;
+			= ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
+			   ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
 		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		      has_attachments = true;
 		    }
@@ -9148,33 +9267,38 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			      break;
 			    if (scp)
 			      continue;
-			    tree d1 = OMP_CLAUSE_DECL (*sc);
-			    tree d2 = OMP_CLAUSE_DECL (c);
-			    while (TREE_CODE (d1) == ARRAY_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    while (TREE_CODE (d2) == ARRAY_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    if (TREE_CODE (d1) == INDIRECT_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    if (TREE_CODE (d2) == INDIRECT_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    while (TREE_CODE (d1) == COMPONENT_REF)
-			      if (TREE_CODE (d2) == COMPONENT_REF
-				  && TREE_OPERAND (d1, 1)
-				     == TREE_OPERAND (d2, 1))
-				{
+			    if ((region_type & ORT_ACC) != 0)
+			      {
+				/* This duplicate checking code is currently only
+				   enabled for OpenACC.  */
+				tree d1 = OMP_CLAUSE_DECL (*sc);
+				tree d2 = OMP_CLAUSE_DECL (c);
+				while (TREE_CODE (d1) == ARRAY_REF)
 				  d1 = TREE_OPERAND (d1, 0);
+				while (TREE_CODE (d2) == ARRAY_REF)
 				  d2 = TREE_OPERAND (d2, 0);
-				}
-			      else
-				break;
-			    if (d1 == d2)
-			      {
-				error_at (OMP_CLAUSE_LOCATION (c),
-					  "%qE appears more than once in map "
-					  "clauses", OMP_CLAUSE_DECL (c));
-				remove = true;
-				break;
+				if (TREE_CODE (d1) == INDIRECT_REF)
+				  d1 = TREE_OPERAND (d1, 0);
+				if (TREE_CODE (d2) == INDIRECT_REF)
+				  d2 = TREE_OPERAND (d2, 0);
+				while (TREE_CODE (d1) == COMPONENT_REF)
+				  if (TREE_CODE (d2) == COMPONENT_REF
+				      && TREE_OPERAND (d1, 1)
+				      == TREE_OPERAND (d2, 1))
+				    {
+				      d1 = TREE_OPERAND (d1, 0);
+				      d2 = TREE_OPERAND (d2, 0);
+				    }
+				  else
+				    break;
+				if (d1 == d2)
+				  {
+				    error_at (OMP_CLAUSE_LOCATION (c),
+					      "%qE appears more than once in map "
+					      "clauses", OMP_CLAUSE_DECL (c));
+				    remove = true;
+				    break;
+				  }
 			      }
 			    if (maybe_lt (offset1, offsetn)
 				|| (known_eq (offset1, offsetn)
@@ -9236,10 +9360,60 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	      break;
 	    }
+	  else
+	    {
+	      /* DECL_P (decl) == true  */
+	      tree *sc;
+	      if (struct_map_to_clause
+		  && (sc = struct_map_to_clause->get (decl)) != NULL
+		  && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
+		  && decl == OMP_CLAUSE_DECL (*sc))
+		{
+		  /* We have found a map of the whole structure after a
+		     leading GOMP_MAP_STRUCT has been created, so refill the
+		     leading clause into a map of the whole structure
+		     variable, and remove the current one.
+		     TODO: we should be able to remove some maps of the
+		     following structure element maps if they are of
+		     compatible TO/FROM/ALLOC type.  */
+		  OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
+		  OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
+		  remove = true;
+		  break;
+		}
+	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
+
+	  if ((code == OMP_TARGET
+	       || code == OMP_TARGET_DATA
+	       || code == OMP_TARGET_ENTER_DATA
+	       || code == OMP_TARGET_EXIT_DATA)
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	    {
+	      for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
+		   octx = octx->outer_context)
+		{
+		  splay_tree_node n
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) OMP_CLAUSE_DECL (c));
+		  /* If this is contained in an outer OpenMP region as a
+		     firstprivate value, remove the attach/detach.  */
+		  if (n && (n->value & GOVD_FIRSTPRIVATE))
+		    {
+		      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      goto do_add;
+		    }
+		}
+
+	      enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
+					     ? GOMP_MAP_DETACH
+					     : GOMP_MAP_ATTACH);
+	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+	    }
+
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 6d0aa8daeb3..c45ee359e60 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -214,6 +214,21 @@ is_oacc_kernels (omp_context *ctx)
 	      == GF_OMP_TARGET_KIND_OACC_KERNELS));
 }
 
+/* Return true if STMT corresponds to an OpenMP target region.  */
+static bool
+is_omp_target (gimple *stmt)
+{
+  if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+    {
+      int kind = gimple_omp_target_kind (stmt);
+      return (kind == GF_OMP_TARGET_KIND_REGION
+	      || kind == GF_OMP_TARGET_KIND_DATA
+	      || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+	      || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+    }
+  return false;
+}
+
 /* If DECL is the artificial dummy VAR_DECL created for non-static
    data member privatization, return the underlying "this" parameter,
    otherwise return NULL.  */
@@ -1346,7 +1361,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && DECL_P (decl)
 	      && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
 		   && (OMP_CLAUSE_MAP_KIND (c)
-		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
 		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
@@ -1367,6 +1384,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && DECL_P (decl)
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt))
+	    {
+	      /* If this is an offloaded region, an attach operation should
+		 only exist when the pointer variable is mapped in a prior
+		 clause.  */
+	      if (is_gimple_omp_offloaded (ctx->stmt))
+		gcc_assert
+		  (maybe_lookup_decl (decl, ctx)
+		   || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+		       && lookup_attribute ("omp declare target",
+					    DECL_ATTRIBUTES (decl))));
+
+	      /* By itself, attach/detach is generated as part of pointer
+		 variable mapping and should not create new variables in the
+		 offloaded region, however sender refs for it must be created
+		 for its address to be passed to the runtime.  */
+	      tree field
+		= build_decl (OMP_CLAUSE_LOCATION (c),
+			      FIELD_DECL, NULL_TREE, ptr_type_node);
+	      SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+	      insert_field_into_struct (ctx->record_type, field);
+	      /* To not clash with a map of the pointer variable itself,
+		 attach/detach maps have their field looked up by the *clause*
+		 tree expression, not the decl.  */
+	      gcc_assert (!splay_tree_lookup (ctx->field_map,
+					      (splay_tree_key) c));
+	      splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+				 (splay_tree_value) field);
+	      break;
+	    }
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
@@ -1606,6 +1657,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable)
 	    break;
+	  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt)
+	      && !is_gimple_omp_offloaded (ctx->stmt))
+	    break;
 	  if (DECL_P (decl))
 	    {
 	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -11458,6 +11514,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_STRUCT:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
 	    break;
 	  case GOMP_MAP_IF_PRESENT:
 	  case GOMP_MAP_FORCE_ALLOC:
@@ -11468,8 +11526,6 @@ 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;
@@ -11524,6 +11580,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    continue;
 	  }
 
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	    && is_omp_target (stmt))
+	  {
+	    gcc_assert (maybe_lookup_field (c, ctx));
+	    map_cnt++;
+	    continue;
+	  }
+
 	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
@@ -11756,14 +11822,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    gcc_assert (DECL_P (ovar2));
 		    ovar = ovar2;
 		  }
-		if (!maybe_lookup_field (ovar, ctx))
+		if (!maybe_lookup_field (ovar, ctx)
+		    && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			 && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+			     || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
 		  continue;
 	      }
 
 	    talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
 	    if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
 	      talign = DECL_ALIGN_UNIT (ovar);
-	    if (nc)
+
+	    if (nc
+		&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		&& is_omp_target (stmt))
+	      {
+		var = lookup_decl_in_outer_ctx (ovar, ctx);
+		x = build_sender_ref (c, ctx);
+		gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+	      }
+	    else if (nc)
 	      {
 		var = lookup_decl_in_outer_ctx (ovar, ctx);
 		x = build_sender_ref (ovar, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 3d64b2e7cb3..679b0505e19 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -21,7 +21,7 @@ void f ()
 
 #pragma acc exit data finalize delete (del_f_p[2:5])
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
-   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
 
 #pragma 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 @@ void f ()
 
 #pragma acc exit data copyout (cpo_f_p[4:10]) finalize
 /* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
-   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
+   { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\)$" 1 "gimple" } } */
 }
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 337c1f7cc77..839269eb62b 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -44,7 +44,7 @@ t1 ()
 }
 
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:\*.*z.? .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:\*.*s\.a.? .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
 /* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
index df405e448b2..9f702ba76f2 100644
--- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -20,8 +20,8 @@ test (int *b, int *c, int *e)
   struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
 
 #pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
-  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*.*s\.b.? \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*.*s\.c.? \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
 
 #pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
-  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
+  /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*.*s\.b.? \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*.*s\.c.? \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index ded1d74ccde..bbc8fb4e32b 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -13,35 +13,35 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
     bar (p);
   #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
     bar (p);
-  #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
+  #pragma omp target map (p) , map (p[0])
     bar (p);
   #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
     bar (&q);
   #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
     bar (p);
-  #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map (t.r)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t.r)
     bar (&t.r);
   #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
   #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
-  #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t)
     bar (t.s);
-  #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map(t.s[0])
     bar (t.s);
   #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
-  #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t.s[2])
     bar (t.s);
-  #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[0:2]) map (t.t[4:6])
     bar (t.t);
-  #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[i:j]) map (t.t[k:l])
     bar (t.t);
   #pragma omp target map (t.s[0]) map (t.r)
     bar (t.s);
@@ -50,5 +50,5 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
   #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
-    bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
+    bar (t.s);
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c
new file mode 100644
index 00000000000..1d9d9252864
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-5.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+  /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses.  */
+  int a, b, c;
+  #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
+  #pragma omp target exit data map(from:a) map(release:b) map(from:c)
+
+  #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
+  a = b = c = 1;
+
+  #pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
+  #pragma omp target exit data map(from:a) map(delete:b) map(from:c)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 373bdcb2114..c5ac06943eb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -12,11 +12,11 @@ program att
 
   !$acc enter data attach(myvar%arr2, myptr)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   !$acc exit data detach(myvar%arr2, myptr)
 ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
 
   ! Test valid usage and processing of the finalize clause.
   !$acc exit data detach(myvar%arr2, myptr) finalize
@@ -24,6 +24,6 @@ program att
   ! For array-descriptor detaches, we no longer generate a "release" mapping
   ! for the pointed-to data for gimplify.c to turn into "delete".  Make sure
   ! the mapping still isn't there.
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
 
 end program att
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index a7788580819..0ff2e471180 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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) 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: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
       END SUBROUTINE f
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
index 73c4f5a87d0..79bab726dea 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
@@ -2,5 +2,5 @@ type t
   integer :: i
 end type t
 type(t) v
-!$omp target enter data map(to:v%i, v%i)  ! { dg-error "appears more than once in map clauses" }
+!$omp target enter data map(to:v%i, v%i)
 end

Re: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Jakub Jelinek]

On Wed, Oct 28, 2020 at 06:32:21PM +0800, Chung-Lin Tang wrote:
> > > @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
> > >   	      /* 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;
> > > -		}
> > So what gimplifies those now?
> 
> They're gimplified somewhere during omp-low now.
> (some gimplify scan testcases were adjusted to accommodate this change)
> 
> I don't remember the exact case I encountered, but there were some issues with gimplified
> expressions inside the map clauses making some later checking more difficult. I haven't seen
> any negative effect of this modification so far.

I don't like that, it goes against many principles, gimplification really
shouldn't leave around non-GIMPLE IL.
If you need to compare same expression or same expression bases later,
perhaps detect the equalities during gimplification before actually gimplifying the
clauses and ensure they are gimplified to the same expression or are using
same base (e.g. by adding SAVE_EXPRs or TARGET_EXPRs before the
gimplification).

	Jakub

Re: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Chung-Lin Tang]

[-- Attachment #1: Type: text/plain, Size: 2593 bytes --]

On 2020/10/29 7:49 PM, Jakub Jelinek wrote:
> On Wed, Oct 28, 2020 at 06:32:21PM +0800, Chung-Lin Tang wrote:
>>>> @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>>>>    	      /* 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;
>>>> -		}
>>> So what gimplifies those now?
>>
>> They're gimplified somewhere during omp-low now.
>> (some gimplify scan testcases were adjusted to accommodate this change)
>>
>> I don't remember the exact case I encountered, but there were some issues with gimplified
>> expressions inside the map clauses making some later checking more difficult. I haven't seen
>> any negative effect of this modification so far.
> 
> I don't like that, it goes against many principles, gimplification really
> shouldn't leave around non-GIMPLE IL.
> If you need to compare same expression or same expression bases later,
> perhaps detect the equalities during gimplification before actually gimplifying the
> clauses and ensure they are gimplified to the same expression or are using
> same base (e.g. by adding SAVE_EXPRs or TARGET_EXPRs before the
> gimplification).

I have moved that same gimplify_expr call down to below the processing block,
and things still work as expected. My aforementioned gimple-scan-test modifications
have all been reverted, and all original tests still pass correctly.

Thanks,
Chung-Lin

	gcc/
	* gimplify.c (is_or_contains_p): New static helper function.
	(omp_target_reorder_clauses): New function.
	(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
	reorder clause list according to OpenMP 5.0 rules. Add handling of
	GOMP_MAP_ATTACH_DETACH for OpenMP cases.
	* omp-low.c (is_omp_target): New static helper function.
	(scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH
	for OpenMP cases.
	(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
	OpenMP cases.

	gcc/testsuite/
	* c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
	* gfortran.dg/gomp/map-2.f90: Likewise.
	* c-c++-common/gomp/map-5.c: New testcase.




[-- Attachment #2: omp5-tgtmapping.02.midend-testsuite.v3.patch --]
[-- Type: text/plain, Size: 23644 bytes --]

diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 29f385c9368..c2500656193 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8364,6 +8364,113 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
   return base;
 }
 
+/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR.  */
+
+static bool
+is_or_contains_p (tree expr, tree base_ptr)
+{
+  while (expr != base_ptr)
+    if (TREE_CODE (base_ptr) == COMPONENT_REF)
+      base_ptr = TREE_OPERAND (base_ptr, 0);
+    else
+      break;
+  return expr == base_ptr;
+}
+
+/* Implement OpenMP 5.x map ordering rules for target directives. There are
+   several rules, and with some level of ambiguity, hopefully we can at least
+   collect the complexity here in one place.  */
+
+static void
+omp_target_reorder_clauses (tree *list_p)
+{
+  /* Collect refs to alloc/release/delete maps.  */
+  auto_vec<tree, 32> ard;
+  tree *cp = list_p;
+  while (*cp != NULL_TREE)
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+	&& (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
+	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
+	    || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
+      {
+	/* Unlink cp and push to ard.  */
+	tree c = *cp;
+	tree nc = OMP_CLAUSE_CHAIN (c);
+	*cp = nc;
+	ard.safe_push (c);
+
+	/* Any associated pointer type maps should also move along.  */
+	while (*cp != NULL_TREE
+	       && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+	       && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
+		   || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
+	  {
+	    c = *cp;
+	    nc = OMP_CLAUSE_CHAIN (c);
+	    *cp = nc;
+	    ard.safe_push (c);
+	  }
+      }
+    else
+      cp = &OMP_CLAUSE_CHAIN (*cp);
+
+  /* Link alloc/release/delete maps to the end of list.  */
+  for (unsigned int i = 0; i < ard.length (); i++)
+    {
+      *cp = ard[i];
+      cp = &OMP_CLAUSE_CHAIN (ard[i]);
+    }
+  *cp = NULL_TREE;
+
+  /* OpenMP 5.0 requires that pointer variables are mapped before
+     its use as a base-pointer.  */
+  auto_vec<tree *, 32> atf;
+  for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+    if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+      {
+	/* Collect alloc, to, from, to/from clause tree pointers.  */
+	gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+	if (k == GOMP_MAP_ALLOC
+	    || k == GOMP_MAP_TO
+	    || k == GOMP_MAP_FROM
+	    || k == GOMP_MAP_TOFROM
+	    || k == GOMP_MAP_ALWAYS_TO
+	    || k == GOMP_MAP_ALWAYS_FROM
+	    || k == GOMP_MAP_ALWAYS_TOFROM)
+	  atf.safe_push (cp);
+      }
+
+  for (unsigned int i = 0; i < atf.length (); i++)
+    if (atf[i])
+      {
+	tree *cp = atf[i];
+	tree decl = OMP_CLAUSE_DECL (*cp);
+	if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
+	  {
+	    tree base_ptr = TREE_OPERAND (decl, 0);
+	    STRIP_TYPE_NOPS (base_ptr);
+	    for (unsigned int j = i + 1; j < atf.length (); j++)
+	      {
+		tree *cp2 = atf[j];
+		tree decl2 = OMP_CLAUSE_DECL (*cp2);
+		if (is_or_contains_p (decl2, base_ptr))
+		  {
+		    /* Move *cp2 to before *cp.  */
+		    tree c = *cp2;
+		    *cp2 = OMP_CLAUSE_CHAIN (c);
+		    OMP_CLAUSE_CHAIN (c) = *cp;
+		    *cp = c;
+		    atf[j] = NULL;
+		  }
+	      }
+	  }
+      }
+}
+
 /* Scan the OMP clauses in *LIST_P, installing mappings into a new
    and previous omp contexts.  */
 
@@ -8405,6 +8512,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	break;
       }
 
+  if (code == OMP_TARGET
+      || code == OMP_TARGET_DATA
+      || code == OMP_TARGET_ENTER_DATA
+      || code == OMP_TARGET_EXIT_DATA)
+    omp_target_reorder_clauses (list_p);
+
   while ((c = *list_p) != NULL)
     {
       bool remove = false;
@@ -8845,15 +8958,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	    }
 	  else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		    || (OMP_CLAUSE_MAP_KIND (c)
-			== GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+			== GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
 		   && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
 	    {
 	      OMP_CLAUSE_SIZE (c)
 		= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
 					   false);
-	      omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
-				GOVD_FIRSTPRIVATE | GOVD_SEEN);
+	      if ((region_type & ORT_TARGET) != 0)
+		omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+				  GOVD_FIRSTPRIVATE | GOVD_SEEN);
 	    }
+
 	  if (!DECL_P (decl))
 	    {
 	      tree d = decl, *pd;
@@ -8878,7 +8994,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      bool indir_p = false;
 	      tree orig_decl = decl;
 	      tree decl_ref = NULL_TREE;
-	      if ((region_type & ORT_ACC) != 0
+	      if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
 		  && TREE_CODE (*pd) == COMPONENT_REF
 		  && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
 		  && code != OACC_UPDATE)
@@ -8886,9 +9002,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  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))
+		      if (((TREE_CODE (decl) == MEM_REF
+			    && integer_zerop (TREE_OPERAND (decl, 1)))
+			   || INDIRECT_REF_P (decl))
+			  && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+			      == POINTER_TYPE))
 			{
 			  indir_p = true;
 			  decl = TREE_OPERAND (decl, 0);
@@ -8915,8 +9033,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      if (decl != orig_decl && DECL_P (decl) && indir_p)
 		{
-		  gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
-							     : GOMP_MAP_ATTACH;
+		  gomp_map_kind k
+		    = ((code == OACC_EXIT_DATA || code == OMP_TARGET_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.  */
@@ -8925,6 +9044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  tree next_clause = OMP_CLAUSE_CHAIN (c);
 		  if (k == GOMP_MAP_ATTACH
 		      && code != OACC_ENTER_DATA
+		      && code != OMP_TARGET_ENTER_DATA
 		      && (!next_clause
 			   || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
 			   || (OMP_CLAUSE_MAP_KIND (next_clause)
@@ -8972,17 +9092,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	      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)
 		  && 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)
+		  && code != OACC_UPDATE
+		  && code != OMP_TARGET_UPDATE)
 		{
 		  if (error_operand_p (decl))
 		    {
@@ -9044,15 +9159,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		  bool has_attachments = false;
 		  /* For OpenACC, pointers in structs should trigger an
 		     attach action.  */
-		  if (attach_detach && (region_type & ORT_ACC) != 0)
+		  if (attach_detach
+		      && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
+			  || code == OMP_TARGET_ENTER_DATA
+			  || code == OMP_TARGET_EXIT_DATA))
+
 		    {
 		      /* 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;
+			= ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
+			   ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
 		      OMP_CLAUSE_SET_MAP_KIND (c, k);
 		      has_attachments = true;
 		    }
@@ -9148,33 +9267,38 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			      break;
 			    if (scp)
 			      continue;
-			    tree d1 = OMP_CLAUSE_DECL (*sc);
-			    tree d2 = OMP_CLAUSE_DECL (c);
-			    while (TREE_CODE (d1) == ARRAY_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    while (TREE_CODE (d2) == ARRAY_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    if (TREE_CODE (d1) == INDIRECT_REF)
-			      d1 = TREE_OPERAND (d1, 0);
-			    if (TREE_CODE (d2) == INDIRECT_REF)
-			      d2 = TREE_OPERAND (d2, 0);
-			    while (TREE_CODE (d1) == COMPONENT_REF)
-			      if (TREE_CODE (d2) == COMPONENT_REF
-				  && TREE_OPERAND (d1, 1)
-				     == TREE_OPERAND (d2, 1))
-				{
+			    if ((region_type & ORT_ACC) != 0)
+			      {
+				/* This duplicate checking code is currently only
+				   enabled for OpenACC.  */
+				tree d1 = OMP_CLAUSE_DECL (*sc);
+				tree d2 = OMP_CLAUSE_DECL (c);
+				while (TREE_CODE (d1) == ARRAY_REF)
 				  d1 = TREE_OPERAND (d1, 0);
+				while (TREE_CODE (d2) == ARRAY_REF)
 				  d2 = TREE_OPERAND (d2, 0);
-				}
-			      else
-				break;
-			    if (d1 == d2)
-			      {
-				error_at (OMP_CLAUSE_LOCATION (c),
-					  "%qE appears more than once in map "
-					  "clauses", OMP_CLAUSE_DECL (c));
-				remove = true;
-				break;
+				if (TREE_CODE (d1) == INDIRECT_REF)
+				  d1 = TREE_OPERAND (d1, 0);
+				if (TREE_CODE (d2) == INDIRECT_REF)
+				  d2 = TREE_OPERAND (d2, 0);
+				while (TREE_CODE (d1) == COMPONENT_REF)
+				  if (TREE_CODE (d2) == COMPONENT_REF
+				      && TREE_OPERAND (d1, 1)
+				      == TREE_OPERAND (d2, 1))
+				    {
+				      d1 = TREE_OPERAND (d1, 0);
+				      d2 = TREE_OPERAND (d2, 0);
+				    }
+				  else
+				    break;
+				if (d1 == d2)
+				  {
+				    error_at (OMP_CLAUSE_LOCATION (c),
+					      "%qE appears more than once in map "
+					      "clauses", OMP_CLAUSE_DECL (c));
+				    remove = true;
+				    break;
+				  }
 			      }
 			    if (maybe_lt (offset1, offsetn)
 				|| (known_eq (offset1, offsetn)
@@ -9220,6 +9344,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 			}
 		    }
 		}
+
+	      if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+		  == GS_ERROR)
+		{
+		  remove = true;
+		  break;
+		}
+
 	      if (!remove
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
 		  && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -9236,10 +9368,60 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 
 	      break;
 	    }
+	  else
+	    {
+	      /* DECL_P (decl) == true  */
+	      tree *sc;
+	      if (struct_map_to_clause
+		  && (sc = struct_map_to_clause->get (decl)) != NULL
+		  && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
+		  && decl == OMP_CLAUSE_DECL (*sc))
+		{
+		  /* We have found a map of the whole structure after a
+		     leading GOMP_MAP_STRUCT has been created, so refill the
+		     leading clause into a map of the whole structure
+		     variable, and remove the current one.
+		     TODO: we should be able to remove some maps of the
+		     following structure element maps if they are of
+		     compatible TO/FROM/ALLOC type.  */
+		  OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
+		  OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
+		  remove = true;
+		  break;
+		}
+	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
 	    flags |= GOVD_MAP_ALWAYS_TO;
+
+	  if ((code == OMP_TARGET
+	       || code == OMP_TARGET_DATA
+	       || code == OMP_TARGET_ENTER_DATA
+	       || code == OMP_TARGET_EXIT_DATA)
+	      && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+	    {
+	      for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
+		   octx = octx->outer_context)
+		{
+		  splay_tree_node n
+		    = splay_tree_lookup (octx->variables,
+					 (splay_tree_key) OMP_CLAUSE_DECL (c));
+		  /* If this is contained in an outer OpenMP region as a
+		     firstprivate value, remove the attach/detach.  */
+		  if (n && (n->value & GOVD_FIRSTPRIVATE))
+		    {
+		      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
+		      goto do_add;
+		    }
+		}
+
+	      enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
+					     ? GOMP_MAP_DETACH
+					     : GOMP_MAP_ATTACH);
+	      OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+	    }
+
 	  goto do_add;
 
 	case OMP_CLAUSE_DEPEND:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 6d0aa8daeb3..c45ee359e60 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -214,6 +214,21 @@ is_oacc_kernels (omp_context *ctx)
 	      == GF_OMP_TARGET_KIND_OACC_KERNELS));
 }
 
+/* Return true if STMT corresponds to an OpenMP target region.  */
+static bool
+is_omp_target (gimple *stmt)
+{
+  if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+    {
+      int kind = gimple_omp_target_kind (stmt);
+      return (kind == GF_OMP_TARGET_KIND_REGION
+	      || kind == GF_OMP_TARGET_KIND_DATA
+	      || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+	      || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+    }
+  return false;
+}
+
 /* If DECL is the artificial dummy VAR_DECL created for non-static
    data member privatization, return the underlying "this" parameter,
    otherwise return NULL.  */
@@ -1346,7 +1361,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && DECL_P (decl)
 	      && ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
 		   && (OMP_CLAUSE_MAP_KIND (c)
-		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+		       != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+		   && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
 		  || TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
 	      && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
@@ -1367,6 +1384,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		  && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
 		break;
 	    }
+	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	      && DECL_P (decl)
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt))
+	    {
+	      /* If this is an offloaded region, an attach operation should
+		 only exist when the pointer variable is mapped in a prior
+		 clause.  */
+	      if (is_gimple_omp_offloaded (ctx->stmt))
+		gcc_assert
+		  (maybe_lookup_decl (decl, ctx)
+		   || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+		       && lookup_attribute ("omp declare target",
+					    DECL_ATTRIBUTES (decl))));
+
+	      /* By itself, attach/detach is generated as part of pointer
+		 variable mapping and should not create new variables in the
+		 offloaded region, however sender refs for it must be created
+		 for its address to be passed to the runtime.  */
+	      tree field
+		= build_decl (OMP_CLAUSE_LOCATION (c),
+			      FIELD_DECL, NULL_TREE, ptr_type_node);
+	      SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+	      insert_field_into_struct (ctx->record_type, field);
+	      /* To not clash with a map of the pointer variable itself,
+		 attach/detach maps have their field looked up by the *clause*
+		 tree expression, not the decl.  */
+	      gcc_assert (!splay_tree_lookup (ctx->field_map,
+					      (splay_tree_key) c));
+	      splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+				 (splay_tree_value) field);
+	      break;
+	    }
 	  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
@@ -1606,6 +1657,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 	      && is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
 	      && varpool_node::get_create (decl)->offloadable)
 	    break;
+	  if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+	       || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	      && is_omp_target (ctx->stmt)
+	      && !is_gimple_omp_offloaded (ctx->stmt))
+	    break;
 	  if (DECL_P (decl))
 	    {
 	      if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -11458,6 +11514,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_STRUCT:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH:
+	  case GOMP_MAP_DETACH:
 	    break;
 	  case GOMP_MAP_IF_PRESENT:
 	  case GOMP_MAP_FORCE_ALLOC:
@@ -11468,8 +11526,6 @@ 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;
@@ -11524,6 +11580,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 	    continue;
 	  }
 
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	    && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+	    && is_omp_target (stmt))
+	  {
+	    gcc_assert (maybe_lookup_field (c, ctx));
+	    map_cnt++;
+	    continue;
+	  }
+
 	if (!maybe_lookup_field (var, ctx))
 	  continue;
 
@@ -11756,14 +11822,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
 		    gcc_assert (DECL_P (ovar2));
 		    ovar = ovar2;
 		  }
-		if (!maybe_lookup_field (ovar, ctx))
+		if (!maybe_lookup_field (ovar, ctx)
+		    && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			 && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+			     || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
 		  continue;
 	      }
 
 	    talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
 	    if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
 	      talign = DECL_ALIGN_UNIT (ovar);
-	    if (nc)
+
+	    if (nc
+		&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		    || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+		&& is_omp_target (stmt))
+	      {
+		var = lookup_decl_in_outer_ctx (ovar, ctx);
+		x = build_sender_ref (c, ctx);
+		gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+	      }
+	    else if (nc)
 	      {
 		var = lookup_decl_in_outer_ctx (ovar, ctx);
 		x = build_sender_ref (ovar, ctx);
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index ded1d74ccde..bbc8fb4e32b 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -13,35 +13,35 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
     bar (p);
   #pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
     bar (p);
-  #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
+  #pragma omp target map (p) , map (p[0])
     bar (p);
   #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
     bar (&q);
   #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
     bar (p);
-  #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map (t.r)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t)
     bar (&t.r);
-  #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.r) map (t.r)
     bar (&t.r);
   #pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
   #pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (&t.r);
-  #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t)
     bar (t.s);
-  #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t) map(t.s[0])
     bar (t.s);
   #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
-  #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.s[0]) map (t.s[2])
     bar (t.s);
-  #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[0:2]) map (t.t[4:6])
     bar (t.t);
-  #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+  #pragma omp target map (t.t[i:j]) map (t.t[k:l])
     bar (t.t);
   #pragma omp target map (t.s[0]) map (t.r)
     bar (t.s);
@@ -50,5 +50,5 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
   #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
     bar (t.s);
   #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
-    bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
+    bar (t.s);
 }
diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c
new file mode 100644
index 00000000000..1d9d9252864
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-5.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+  /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses.  */
+  int a, b, c;
+  #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
+  #pragma omp target exit data map(from:a) map(release:b) map(from:c)
+
+  #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
+  a = b = c = 1;
+
+  #pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
+  #pragma omp target exit data map(from:a) map(delete:b) map(from:c)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
index 73c4f5a87d0..79bab726dea 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
@@ -2,5 +2,5 @@ type t
   integer :: i
 end type t
 type(t) v
-!$omp target enter data map(to:v%i, v%i)  ! { dg-error "appears more than once in map clauses" }
+!$omp target enter data map(to:v%i, v%i)
 end

Re: [PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases

[Jakub Jelinek]

On Wed, Nov 04, 2020 at 02:02:56AM +0800, Chung-Lin Tang wrote:
> 	gcc/
> 	* gimplify.c (is_or_contains_p): New static helper function.
> 	(omp_target_reorder_clauses): New function.
> 	(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
> 	reorder clause list according to OpenMP 5.0 rules. Add handling of
> 	GOMP_MAP_ATTACH_DETACH for OpenMP cases.
> 	* omp-low.c (is_omp_target): New static helper function.
> 	(scan_sharing_clauses): Add scan phase handling of GOMP_MAP_ATTACH/DETACH
> 	for OpenMP cases.
> 	(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
> 	OpenMP cases.
> 
> 	gcc/testsuite/
> 	* c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
> 	* gfortran.dg/gomp/map-2.f90: Likewise.
> 	* c-c++-common/gomp/map-5.c: New testcase.

Ok, thanks.

	Jakub