openmp: Implement OpenMP 5.0 base-pointer attachement and clause ordering

This patch implements some parts of the target variable mapping changes
specified in OpenMP 5.0, including base-pointer attachment/detachment
behavior for array section list-items in map clauses, and ordering of
map clauses according to map kind.

2020-11-10  Chung-Lin Tang  <cltang@codesourcery.com>

gcc/c-family/ChangeLog:

	* c-common.h (c_omp_adjust_map_clauses): New declaration.
	* c-omp.c (struct map_clause): Helper type for c_omp_adjust_map_clauses.
	(c_omp_adjust_map_clauses): New function.

gcc/c/ChangeLog:

	* c-parser.c (c_parser_omp_target_data): Add use of
	new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as
	handled map clause kind.
	(c_parser_omp_target_enter_data): Likewise.
	(c_parser_omp_target_exit_data): Likewise.
	(c_parser_omp_target): Likewise.
	* c-typeck.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
	use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type.
	(c_finish_omp_clauses): Adjust bitmap checks to allow struct decl and
	same struct field access to co-exist on OpenMP construct.

gcc/cp/ChangeLog:

	* parser.c (cp_parser_omp_target_data): Add use of
	new c_omp_adjust_map_clauses function. Add GOMP_MAP_ATTACH_DETACH as
	handled map clause kind.
	(cp_parser_omp_target_enter_data): Likewise.
	(cp_parser_omp_target_exit_data): Likewise.
	(cp_parser_omp_target): Likewise.
	* semantics.c (handle_omp_array_sections): Adjust COMPONENT_REF case to
	use GOMP_MAP_ATTACH_DETACH map kind for C_ORT_OMP region type. Fix
	interaction between reference case and attach/detach.
	(finish_omp_clauses): Adjust bitmap checks to allow struct decl and
	same struct field access to co-exist on OpenMP construct.

gcc/ChangeLog:

	* 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/ChangeLog:

	* 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.

libgomp/ChangeLog:

	* libgomp.h (enum gomp_map_vars_kind): Adjust enum values to be bit-flag
	usable.
	* oacc-mem.c (acc_map_data): Adjust gomp_map_vars argument flags to
	'GOMP_MAP_VARS_OPENACC | GOMP_MAP_VARS_ENTER_DATA'.
	(goacc_enter_datum): Likewise for call to gomp_map_vars_async.
	(goacc_enter_data_internal): Likewise.
	* target.c (gomp_map_vars_internal):
	Change checks of GOMP_MAP_VARS_ENTER_DATA to use bit-and (&). Adjust use
	of gomp_attach_pointer for OpenMP cases.
	(gomp_exit_data): Add handling of GOMP_MAP_DETACH.
	(GOMP_target_enter_exit_data): Add handling of GOMP_MAP_ATTACH.
	* testsuite/libgomp.c-c++-common/ptr-attach-1.c: New testcase.

11 files changed, 493 insertions, 89 deletions

diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index bfcc279..b80db23 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1224,6 +1224,7 @@ extern enum omp_clause_defaultmap_kind c_omp_predetermined_mapping (tree);
 extern tree c_omp_check_context_selector (location_t, tree);
 extern void c_omp_mark_declare_variant (location_t, tree, tree);
 extern const char *c_omp_map_clause_name (tree, bool);
+extern void c_omp_adjust_map_clauses (tree, bool);
 
 /* Return next tree in the chain for chain_next walking of tree nodes.  */
 static inline tree
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index da5564b..8457211 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2771,3 +2771,93 @@ c_omp_map_clause_name (tree clause, bool oacc)
     }
   return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
 }
+
+/* Used to merge map clause information in c_omp_adjust_map_clauses.  */
+struct map_clause
+{
+  tree clause;
+  bool firstprivate_ptr_p;
+  bool decl_mapped;
+  bool omp_declare_target;
+  map_clause (void) : clause (NULL_TREE), firstprivate_ptr_p (false),
+    decl_mapped (false), omp_declare_target (false) { }
+};
+
+/* Adjust map clauses after normal clause parsing, mainly to turn specific
+   base-pointer map cases into attach/detach and mark them addressable.  */
+void
+c_omp_adjust_map_clauses (tree clauses, bool is_target)
+{
+  if (!is_target)
+    {
+      /* If this is not a target construct, just turn firstprivate pointers
+	 into attach/detach, the runtime will check and do the rest.  */
+
+      for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+	if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	    && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	    && DECL_P (OMP_CLAUSE_DECL (c))
+	    && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+	  {
+	    tree ptr = OMP_CLAUSE_DECL (c);
+	    OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ATTACH_DETACH);
+	    c_common_mark_addressable_vec (ptr);
+	  }
+      return;
+    }
+
+  hash_map<tree, map_clause> maps;
+
+  for (tree c = clauses; c; c = OMP_CLAUSE_CHAIN (c))
+    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+	&& DECL_P (OMP_CLAUSE_DECL (c)))
+      {
+	/* If this is for a target construct, the firstprivate pointer
+	   is changed to attach/detach if either is true:
+	   (1) the base-pointer is mapped in this same construct, or
+	   (2) the base-pointer is a variable place on the device by
+	       "declare target" directives.
+
+	   Here we iterate through all map clauses collecting these cases,
+	   and merge them with a hash_map to process below.  */
+
+	if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+	    && POINTER_TYPE_P (TREE_TYPE (OMP_CLAUSE_DECL (c))))
+	  {
+	    tree ptr = OMP_CLAUSE_DECL (c);
+	    map_clause &mc = maps.get_or_insert (ptr);
+	    if (mc.clause == NULL_TREE)
+	      mc.clause = c;
+	    mc.firstprivate_ptr_p = true;
+
+	    if (is_global_var (ptr)
+		&& lookup_attribute ("omp declare target",
+				     DECL_ATTRIBUTES (ptr)))
+	      mc.omp_declare_target = true;
+	  }
+	else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALLOC
+		 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO
+		 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FROM
+		 || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TOFROM
+		 || 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)
+	  {
+	    map_clause &mc = maps.get_or_insert (OMP_CLAUSE_DECL (c));
+	    mc.decl_mapped = true;
+	  }
+      }
+
+  for (hash_map<tree, map_clause>::iterator i = maps.begin ();
+       i != maps.end (); ++i)
+    {
+      map_clause &mc = (*i).second;
+
+      if (mc.firstprivate_ptr_p
+	  && (mc.decl_mapped || mc.omp_declare_target))
+	{
+	  OMP_CLAUSE_SET_MAP_KIND (mc.clause, GOMP_MAP_ATTACH_DETACH);
+	  c_common_mark_addressable_vec (OMP_CLAUSE_DECL (mc.clause));
+	}
+    }
+}
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ecc3d21..377914c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -19511,6 +19511,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				"#pragma omp target data");
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19528,6 +19529,7 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p)
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19651,6 +19653,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
 				"#pragma omp target enter data");
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19664,6 +19667,7 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser,
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19735,7 +19739,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
   tree clauses
     = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				"#pragma omp target exit data");
-
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -19750,6 +19754,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser,
 	    break;
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -19960,6 +19965,8 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p)
   OMP_TARGET_CLAUSES (stmt)
     = c_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				"#pragma omp target");
+  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
   pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level ();
   block = c_begin_compound_stmt (true);
@@ -19984,6 +19991,7 @@ check_clauses:
 	  case GOMP_MAP_ALLOC:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 9684037..df1dad4 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13584,11 +13584,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
       if (ort != C_ORT_OMP && ort != C_ORT_ACC)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
       else if (TREE_CODE (t) == COMPONENT_REF)
-	{
-	  gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
-					       : GOMP_MAP_ALWAYS_POINTER;
-	  OMP_CLAUSE_SET_MAP_KIND (c2, k);
-	}
+	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -14711,7 +14707,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		break;
 	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 		{
-		  if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+		  if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+		      || (ort == C_ORT_OMP
+			  && bitmap_bit_p (&map_head, DECL_UID (t))))
 		    break;
 		}
 	    }
@@ -14780,7 +14778,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      else
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
-	  else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
+		   && (ort != C_ORT_OMP
+		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -14794,7 +14794,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	    }
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+		   && ort == C_ORT_ACC)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qD appears more than once in data clauses", t);
+	      remove = true;
+	    }
+	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 4c819ea..9c08c0e 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -40785,6 +40785,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK,
 				 "#pragma omp target data", pragma_tok);
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -40803,6 +40804,7 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p)
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -40886,6 +40888,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK,
 				 "#pragma omp target enter data", pragma_tok);
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -40900,6 +40903,7 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok,
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -40974,6 +40978,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
   tree clauses
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK,
 				 "#pragma omp target exit data", pragma_tok);
+  c_omp_adjust_map_clauses (clauses, false);
   int map_seen = 0;
   for (tree *pc = &clauses; *pc;)
     {
@@ -40989,6 +40994,7 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok,
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    map_seen |= 1;
@@ -41238,6 +41244,8 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
   OMP_TARGET_CLAUSES (stmt)
     = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
 				 "#pragma omp target", pragma_tok);
+  c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
+
   pc = &OMP_TARGET_CLAUSES (stmt);
   keep_next_level (true);
   OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
@@ -41261,6 +41269,7 @@ check_clauses:
 	  case GOMP_MAP_FIRSTPRIVATE_POINTER:
 	  case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
 	  case GOMP_MAP_ALWAYS_POINTER:
+	  case GOMP_MAP_ATTACH_DETACH:
 	    break;
 	  default:
 	    error_at (OMP_CLAUSE_LOCATION (*pc),
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index a550db6..33d715e 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5382,11 +5382,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	  if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
 	  else if (TREE_CODE (t) == COMPONENT_REF)
-	    {
-	      gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
-						   : GOMP_MAP_ALWAYS_POINTER;
-	      OMP_CLAUSE_SET_MAP_KIND (c2, k);
-	    }
+	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 	  else if (REFERENCE_REF_P (t)
 		   && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
 	    {
@@ -5424,8 +5420,12 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 					  OMP_CLAUSE_MAP);
 	      OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
 	      OMP_CLAUSE_DECL (c3) = ptr;
-	      if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER)
-		OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+	      if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER
+		  || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH)
+		{
+		  OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+		  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+		}
 	      else
 		OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
 	      OMP_CLAUSE_SIZE (c3) = size_zero_node;
@@ -7486,7 +7486,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      t = TREE_OPERAND (t, 0);
 	      OMP_CLAUSE_DECL (c) = t;
 	    }
-	  if (ort == C_ORT_ACC
+	  if ((ort == C_ORT_ACC || ort == C_ORT_OMP)
 	      && TREE_CODE (t) == COMPONENT_REF
 	      && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
 	    t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
@@ -7532,7 +7532,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		t = TREE_OPERAND (t, 0);
 	      if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
 		{
-		  if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+		  if (bitmap_bit_p (&map_field_head, DECL_UID (t))
+		      || (ort == C_ORT_OMP
+			  && bitmap_bit_p (&map_head, DECL_UID (t))))
 		    goto handle_map_references;
 		}
 	    }
@@ -7626,13 +7628,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		bitmap_set_bit (&generic_head, DECL_UID (t));
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
-		   && (ort != C_ORT_ACC
-		       || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
+		   && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
 	    {
 	      if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in motion clauses", t);
-	      if (ort == C_ORT_ACC)
+	      else if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),
 			  "%qD appears more than once in data clauses", t);
 	      else
@@ -7641,7 +7642,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	      remove = true;
 	    }
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
-		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+		   && ort == C_ORT_ACC)
+	    {
+	      error_at (OMP_CLAUSE_LOCATION (c),
+			"%qD appears more than once in data clauses", t);
+	      remove = true;
+	    }
+	  else if (bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
 	      if (ort == C_ORT_ACC)
 		error_at (OMP_CLAUSE_LOCATION (c),
@@ -7677,17 +7684,14 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		       && (OMP_CLAUSE_MAP_KIND (c)
 			   != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
 		       && (OMP_CLAUSE_MAP_KIND (c)
-			   != GOMP_MAP_ALWAYS_POINTER))
+			   != GOMP_MAP_ALWAYS_POINTER)
+		       && (OMP_CLAUSE_MAP_KIND (c)
+			   != GOMP_MAP_ATTACH_DETACH))
 		{
 		  tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 					      OMP_CLAUSE_MAP);
 		  if (TREE_CODE (t) == COMPONENT_REF)
-		    {
-		      gomp_map_kind k
-			= (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
-					     : GOMP_MAP_ALWAYS_POINTER;
-		      OMP_CLAUSE_SET_MAP_KIND (c2, k);
-		    }
+		    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
 		  else
 		    OMP_CLAUSE_SET_MAP_KIND (c2,
 					     GOMP_MAP_FIRSTPRIVATE_REFERENCE);
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index aa3b914..b2c623b 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 ea9008b..447d7db 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
@@ -1368,6 +1385,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
 		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)))
@@ -1607,6 +1658,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
@@ -11471,6 +11527,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:
@@ -11481,8 +11539,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;
@@ -11537,6 +11593,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;
 
@@ -11769,14 +11835,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 ded1d74..bbc8fb4 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 0000000..1d9d925
--- /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 73c4f5a..79bab72 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