diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index 1b4a1124c25..28fbb1d0537 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -1989,6 +1989,16 @@ c_omp_split_clauses (location_t loc, enum tree_code code,
 			"%<parallel for%>, %<parallel for simd%>");
 	      OMP_CLAUSE_REDUCTION_INSCAN (clauses) = 0;
 	    }
+	  if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_MAP)) != 0)
+	    {
+	      c = build_omp_clause (OMP_CLAUSE_LOCATION (clauses),
+				    OMP_CLAUSE_MAP);
+	      OMP_CLAUSE_DECL (c) = OMP_CLAUSE_DECL (clauses);
+	      OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+	      OMP_CLAUSE_MAP_IMPLICIT (c) = 1;
+	      OMP_CLAUSE_CHAIN (c) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
+	      cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = c;
+	    }
 	  if ((mask & (OMP_CLAUSE_MASK_1 << PRAGMA_OMP_CLAUSE_SCHEDULE)) != 0)
 	    {
 	      if (code == OMP_SIMD)
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index bee01df1234..5f322874423 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -13657,6 +13657,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ATTACH_DETACH);
       else
 	OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+      OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
       if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	  && !c_mark_addressable (t))
 	return false;
@@ -13927,7 +13928,8 @@ tree
 c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
-  bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
+  bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
+  bitmap_head oacc_reduction_head;
   tree c, t, type, *pc;
   tree simdlen = NULL_TREE, safelen = NULL_TREE;
   bool branch_seen = false;
@@ -13947,7 +13949,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
      has been seen, -2 if mixed inscan/normal reduction diagnosed.  */
   int reduction_seen = 0;
   bool allocate_seen = false;
-  bool firstprivate_implicit_moved = false;
+  bool implicit_moved = false;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -13957,6 +13959,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead.  */
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+  bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack);
   /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
      instead.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
@@ -14389,17 +14392,25 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  break;
 
 	case OMP_CLAUSE_FIRSTPRIVATE:
-	  if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
-	      && !firstprivate_implicit_moved)
+	  if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved)
 	    {
-	      firstprivate_implicit_moved = true;
-	      /* Move firstprivate clauses with
-		 OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of
+	    move_implicit:
+	      implicit_moved = true;
+	      /* Move firstprivate and map clauses with
+		 OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of
 		 clauses chain.  */
-	      tree cl = NULL, *pc1 = pc, *pc2 = &cl;
+	      tree cl1 = NULL_TREE, cl2 = NULL_TREE;
+	      tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2;
 	      while (*pc1)
 		if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE
 		    && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1))
+		  {
+		    *pc3 = *pc1;
+		    pc3 = &OMP_CLAUSE_CHAIN (*pc3);
+		    *pc1 = OMP_CLAUSE_CHAIN (*pc1);
+		  }
+		else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP
+			 && OMP_CLAUSE_MAP_IMPLICIT (*pc1))
 		  {
 		    *pc2 = *pc1;
 		    pc2 = &OMP_CLAUSE_CHAIN (*pc2);
@@ -14407,10 +14418,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  }
 		else
 		  pc1 = &OMP_CLAUSE_CHAIN (*pc1);
-	      *pc2 = NULL;
-	      *pc1 = cl;
-	      if (pc1 != pc)
-		continue;
+	      *pc3 = NULL;
+	      *pc2 = cl2;
+	      *pc1 = cl1;
+	      continue;
 	    }
 	  t = OMP_CLAUSE_DECL (c);
 	  need_complete = true;
@@ -14421,6 +14432,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qE is not a variable in clause %<firstprivate%>", t);
 	      remove = true;
 	    }
+	  else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
+		   && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c)
+		   && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
+	    remove = true;
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
 		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
@@ -14673,6 +14688,9 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  break;
 
 	case OMP_CLAUSE_MAP:
+	  if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
+	    goto move_implicit;
+	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE__CACHE_:
@@ -14706,6 +14724,16 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		    {
 		      while (TREE_CODE (t) == COMPONENT_REF)
 			t = TREE_OPERAND (t, 0);
+		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			  && OMP_CLAUSE_MAP_IMPLICIT (c)
+			  && (bitmap_bit_p (&map_head, DECL_UID (t))
+			      || bitmap_bit_p (&map_field_head, DECL_UID (t))
+			      || bitmap_bit_p (&map_firstprivate_head,
+					       DECL_UID (t))))
+			{
+			  remove = true;
+			  break;
+			}
 		      if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -14859,6 +14887,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
 	      remove = true;
 	    }
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		   && OMP_CLAUSE_MAP_IMPLICIT (c)
+		   && (bitmap_bit_p (&map_head, DECL_UID (t))
+		       || bitmap_bit_p (&map_field_head, DECL_UID (t))
+		       || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))))
+	    remove = true;
 	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
@@ -14880,7 +14914,10 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  remove = true;
 		}
 	      else
-		bitmap_set_bit (&generic_head, DECL_UID (t));
+		{
+		  bitmap_set_bit (&generic_head, DECL_UID (t));
+		  bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
+		}
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		   && (ort != C_ORT_OMP
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 6fafd0648e0..fe370a21366 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -5548,6 +5548,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	    }
 	  else
 	    OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+	  OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
 	  if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
 	      && !cxx_mark_addressable (t))
 	    return false;
@@ -5574,6 +5575,7 @@ handle_omp_array_sections (tree c, enum c_omp_region_type ort)
 	      tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
 					  OMP_CLAUSE_MAP);
 	      OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
+	      OMP_CLAUSE_MAP_IMPLICIT (c2) = OMP_CLAUSE_MAP_IMPLICIT (c);
 	      OMP_CLAUSE_DECL (c3) = ptr;
 	      if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER
 		  || OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ATTACH_DETACH)
@@ -6510,7 +6512,8 @@ tree
 finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 {
   bitmap_head generic_head, firstprivate_head, lastprivate_head;
-  bitmap_head aligned_head, map_head, map_field_head, oacc_reduction_head;
+  bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head;
+  bitmap_head oacc_reduction_head;
   tree c, t, *pc;
   tree safelen = NULL_TREE;
   bool branch_seen = false;
@@ -6527,7 +6530,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   bool allocate_seen = false;
   tree detach_seen = NULL_TREE;
   bool mergeable_seen = false;
-  bool firstprivate_implicit_moved = false;
+  bool implicit_moved = false;
 
   bitmap_obstack_initialize (NULL);
   bitmap_initialize (&generic_head, &bitmap_default_obstack);
@@ -6537,6 +6540,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
   /* If ort == C_ORT_OMP_DECLARE_SIMD used as uniform_head instead.  */
   bitmap_initialize (&map_head, &bitmap_default_obstack);
   bitmap_initialize (&map_field_head, &bitmap_default_obstack);
+  bitmap_initialize (&map_firstprivate_head, &bitmap_default_obstack);
   /* If ort == C_ORT_OMP used as nontemporal_head or use_device_xxx_head
      instead.  */
   bitmap_initialize (&oacc_reduction_head, &bitmap_default_obstack);
@@ -6852,17 +6856,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  break;
 
 	case OMP_CLAUSE_FIRSTPRIVATE:
-	  if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
-	      && !firstprivate_implicit_moved)
+	  if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c) && !implicit_moved)
 	    {
-	      firstprivate_implicit_moved = true;
-	      /* Move firstprivate clauses with
-		 OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT set to the end of
+	    move_implicit:
+	      implicit_moved = true;
+	      /* Move firstprivate and map clauses with
+		 OMP_CLAUSE_{FIRSTPRIVATE,MAP}_IMPLICIT set to the end of
 		 clauses chain.  */
-	      tree cl = NULL, *pc1 = pc, *pc2 = &cl;
+	      tree cl1 = NULL_TREE, cl2 = NULL_TREE;
+	      tree *pc1 = pc, *pc2 = &cl1, *pc3 = &cl2;
 	      while (*pc1)
 		if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_FIRSTPRIVATE
 		    && OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (*pc1))
+		  {
+		    *pc3 = *pc1;
+		    pc3 = &OMP_CLAUSE_CHAIN (*pc3);
+		    *pc1 = OMP_CLAUSE_CHAIN (*pc1);
+		  }
+		else if (OMP_CLAUSE_CODE (*pc1) == OMP_CLAUSE_MAP
+			 && OMP_CLAUSE_MAP_IMPLICIT (*pc1))
 		  {
 		    *pc2 = *pc1;
 		    pc2 = &OMP_CLAUSE_CHAIN (*pc2);
@@ -6870,10 +6882,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  }
 		else
 		  pc1 = &OMP_CLAUSE_CHAIN (*pc1);
-	      *pc2 = NULL;
-	      *pc1 = cl;
-	      if (pc1 != pc)
-		continue;
+	      *pc3 = NULL;
+	      *pc2 = cl2;
+	      *pc1 = cl1;
+	      continue;
 	    }
 	  t = omp_clause_decl_field (OMP_CLAUSE_DECL (c));
 	  if (t)
@@ -6904,6 +6916,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			  t);
 	      remove = true;
 	    }
+	  else if (OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (c)
+		   && !OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT_TARGET (c)
+		   && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t)))
+	    remove = true;
 	  else if (bitmap_bit_p (&generic_head, DECL_UID (t))
 		   || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
 	    {
@@ -7620,6 +7636,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 	  break;
 
 	case OMP_CLAUSE_MAP:
+	  if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
+	    goto move_implicit;
+	  /* FALLTHRU */
 	case OMP_CLAUSE_TO:
 	case OMP_CLAUSE_FROM:
 	case OMP_CLAUSE__CACHE_:
@@ -7651,6 +7670,16 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			t = TREE_OPERAND (t, 0);
 		      if (REFERENCE_REF_P (t))
 			t = TREE_OPERAND (t, 0);
+		      if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			  && OMP_CLAUSE_MAP_IMPLICIT (c)
+			  && (bitmap_bit_p (&map_head, DECL_UID (t))
+			      || bitmap_bit_p (&map_field_head, DECL_UID (t))
+			      || bitmap_bit_p (&map_firstprivate_head,
+					       DECL_UID (t))))
+			{
+			  remove = true;
+			  break;
+			}
 		      if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
 			break;
 		      if (bitmap_bit_p (&map_head, DECL_UID (t)))
@@ -7831,6 +7860,13 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 			"%qD is not a pointer variable", t);
 	      remove = true;
 	    }
+	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+		   && OMP_CLAUSE_MAP_IMPLICIT (c)
+		   && (bitmap_bit_p (&map_head, DECL_UID (t))
+		       || bitmap_bit_p (&map_field_head, DECL_UID (t))
+		       || bitmap_bit_p (&map_firstprivate_head,
+					DECL_UID (t))))
+	    remove = true;
 	  else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
 		   && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
 	    {
@@ -7852,7 +7888,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
 		  remove = true;
 		}
 	      else
-		bitmap_set_bit (&generic_head, DECL_UID (t));
+		{
+		  bitmap_set_bit (&generic_head, DECL_UID (t));
+		  bitmap_set_bit (&map_firstprivate_head, DECL_UID (t));
+		}
 	    }
 	  else if (bitmap_bit_p (&map_head, DECL_UID (t))
 		   && !bitmap_bit_p (&map_field_head, DECL_UID (t)))
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
index d2987a1d9b4..6c44600b07e 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-10.c
@@ -95,22 +95,22 @@ bar (void)
     #pragma omp section
     r12[1]++;
   }
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 60\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r13 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r13 \\+ 4" "gimple" } } */
   #pragma omp target parallel reduction(+:r13[1:15])
   r13[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 64\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r14 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r14 \\+ 4" "gimple" } } *//* FIXME.  */
   #pragma omp target parallel for reduction(+:r14[1:16])
   for (int i = 0; i < 64; i++)
     r14[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 68\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r15 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r15 \\+ 4" "gimple" } } *//* FIXME.  */
@@ -118,31 +118,31 @@ bar (void)
   #pragma omp target parallel for simd reduction(+:r15[1:17])
   for (int i = 0; i < 64; i++)
     r15[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 72\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r16 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization.  */
   /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail.  */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail.  */
   #pragma omp target parallel loop reduction(+:r16[1:18])
   for (int i = 0; i < 64; i++)
     r16[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 76\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r17 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r17 \\+ 4" "gimple" } } */
   #pragma omp target teams reduction(+:r17[1:19])
   r17[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 80\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r18 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r18 \\+ 4" "gimple" } } */
   #pragma omp target teams distribute reduction(+:r18[1:20])
   for (int i = 0; i < 64; i++)
     r18[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 84\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r19 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r19 \\+ 4" "gimple" } } */
@@ -151,8 +151,8 @@ bar (void)
   #pragma omp target teams distribute parallel for reduction(+:r19[1:21])
   for (int i = 0; i < 64; i++)
     r19[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 88\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r20 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r20 \\+ 4" "gimple" } } */
@@ -162,8 +162,8 @@ bar (void)
   #pragma omp target teams distribute parallel for simd reduction(+:r20[1:22])
   for (int i = 0; i < 64; i++)
     r20[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 92\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r21 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r21 \\+ 4" "gimple" } } */
@@ -171,9 +171,9 @@ bar (void)
   #pragma omp target teams distribute simd reduction(+:r21[1:23])
   for (int i = 0; i < 64; i++)
     r21[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 96\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r22 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* FIXME: Should be shared, but firstprivate is an optimization.  */
   /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail.  */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail.  */
@@ -182,8 +182,8 @@ bar (void)
   #pragma omp target teams loop reduction(+:r22[1:24])
   for (int i = 0; i < 64; i++)
     r22[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:\\\*_\[0-9]* \\\[len: 100\\\]" "gimple" } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(firstprivate:r23 \\\[pointer assign, bias: 4\\\]\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*\\)r23 \\+ 4" "gimple" } } */
   #pragma omp target simd reduction(+:r23[1:25])
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-16.c b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c
new file mode 100644
index 00000000000..84cd85da7cf
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-16.c
@@ -0,0 +1,16 @@
+/* PR middle-end/99928 */
+
+void
+foo (void)
+{
+  int a[6] = {};
+  #pragma omp target simd reduction(+:a[:3])
+  for (int i = 0; i < 6; i++)
+    a[0]++;
+  #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a)
+  for (int i = 0; i < 6; i++)
+    a[0]++;
+  #pragma omp target simd reduction(+:a[:3]) map(always, tofrom: a[:6])
+  for (int i = 0; i < 6; i++)
+    a[0]++;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
index fc57573eaf5..27e6ad1eb36 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-8.c
@@ -94,48 +94,48 @@ bar (void)
     #pragma omp section
     r12++;
   }
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r13\\)" "gimple" } } */
   #pragma omp target parallel reduction(+:r13)
   r13++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r14\\)" "gimple" } } *//* FIXME.  */
   #pragma omp target parallel for reduction(+:r14)
   for (int i = 0; i < 64; i++)
     r14++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } *//* FIXME.  */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r15\\)" "gimple" } } */
   #pragma omp target parallel for simd reduction(+:r15)
   for (int i = 0; i < 64; i++)
     r15++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail.  */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r16\\)" "gimple" } } *//* NOTE: This is implementation detail.  */
   #pragma omp target parallel loop reduction(+:r16)
   for (int i = 0; i < 64; i++)
     r16++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r17\\)" "gimple" } } */
   #pragma omp target teams reduction(+:r17)
   r17++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r18\\)" "gimple" } } */
   #pragma omp target teams distribute reduction(+:r18)
   for (int i = 0; i < 64; i++)
     r18++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r19\\)" "gimple" } } *//* FIXME: This should be on for instead.  */
@@ -143,8 +143,8 @@ bar (void)
   #pragma omp target teams distribute parallel for reduction(+:r19)
   for (int i = 0; i < 64; i++)
     r19++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:r20\\)" "gimple" } } *//* FIXME: This should be on for instead.  */
@@ -153,16 +153,16 @@ bar (void)
   #pragma omp target teams distribute parallel for simd reduction(+:r20)
   for (int i = 0; i < 64; i++)
     r20++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r21\\)" "gimple" } } */
   #pragma omp target teams distribute simd reduction(+:r21)
   for (int i = 0; i < 64; i++)
     r21++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:r22\\)" "gimple" } } *//* NOTE: This is implementation detail.  */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r22\\)" "gimple" } } *//* NOTE: This is implementation detail.  */
@@ -171,8 +171,8 @@ bar (void)
   #pragma omp target teams loop reduction(+:r22)
   for (int i = 0; i < 64; i++)
     r22++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" { xfail *-*-* } } } */
-  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23" "gimple" } } */
+  /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:r23\\)" "gimple" } } */
   #pragma omp target simd reduction(+:r23)
   for (int i = 0; i < 64; i++)
diff --git a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
index c049bb0d257..86235277667 100644
--- a/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
+++ b/gcc/testsuite/c-c++-common/gomp/pr99928-9.c
@@ -94,19 +94,19 @@ bar (void)
     #pragma omp section
     r12[1]++;
   }
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r13\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r13\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r13 \\+ 4" "gimple" } } */
   #pragma omp target parallel reduction(+:r13[1:2])
   r13[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r14\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r14" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r14 \\+ 4" "gimple" } } *//* FIXME.  */
   #pragma omp target parallel for reduction(+:r14[1:2])
   for (int i = 0; i < 64; i++)
     r14[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r15\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r15\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME: This should be on for instead.  */
   /* { dg-final { scan-tree-dump-not "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r15 \\+ 4" "gimple" } } *//* FIXME.  */
@@ -114,7 +114,7 @@ bar (void)
   #pragma omp target parallel for simd reduction(+:r15[1:2])
   for (int i = 0; i < 64; i++)
     r15[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r16\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r16\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp parallel\[^\n\r]*shared\\(r16\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp for\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r16 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail.  */
@@ -122,19 +122,19 @@ bar (void)
   #pragma omp target parallel loop reduction(+:r16[1:2])
   for (int i = 0; i < 64; i++)
     r16[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r17\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r17\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r17 \\+ 4" "gimple" } } */
   #pragma omp target teams reduction(+:r17[1:2])
   r17[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r18\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r18\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r18 \\+ 4" "gimple" } } */
   #pragma omp target teams distribute reduction(+:r18[1:2])
   for (int i = 0; i < 64; i++)
     r18[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r19\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r19\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r19 \\+ 4" "gimple" } } */
@@ -143,7 +143,7 @@ bar (void)
   #pragma omp target teams distribute parallel for reduction(+:r19[1:2])
   for (int i = 0; i < 64; i++)
     r19[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r20\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r20\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r20 \\+ 4" "gimple" } } */
@@ -153,7 +153,7 @@ bar (void)
   #pragma omp target teams distribute parallel for simd reduction(+:r20[1:2])
   for (int i = 0; i < 64; i++)
     r20[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r21\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r21\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r21 \\+ 4" "gimple" } } */
@@ -161,7 +161,7 @@ bar (void)
   #pragma omp target teams distribute simd reduction(+:r21[1:2])
   for (int i = 0; i < 64; i++)
     r21[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r22\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r22\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp teams\[^\n\r]*shared\\(r22\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp distribute\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r22 \\+ 4" "gimple" } } *//* NOTE: This is implementation detail.  */
@@ -171,7 +171,7 @@ bar (void)
   #pragma omp target teams loop reduction(+:r22[1:2])
   for (int i = 0; i < 64; i++)
     r22[1]++;
-  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" { xfail *-*-* } } } */
+  /* { dg-final { scan-tree-dump "omp target\[^\n\r]*map\\(tofrom:r23\\\[1\\\] \\\[len: 8\\\]" "gimple" } } */
   /* { dg-final { scan-tree-dump-not "omp target\[^\n\r]*firstprivate\\(r23\\)" "gimple" } } */
   /* { dg-final { scan-tree-dump "omp simd\[^\n\r]*reduction\\(\\+:MEM\[^\n\r]*&r23 \\+ 4" "gimple" } } */
   #pragma omp target simd reduction(+:r23[1:2])
diff --git a/gcc/tree.h b/gcc/tree.h
index 37aca8963fe..5935d0e9f7c 100644
--- a/gcc/tree.h
+++ b/gcc/tree.h
@@ -1654,6 +1654,11 @@ class auto_suppress_location_wrappers
    variable.  */
 #define OMP_CLAUSE_MAP_IN_REDUCTION(NODE) \
   TREE_PRIVATE (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP))
+/* Nonzero on map clauses added implicitly for reduction clauses on combined
+   or composite constructs.  They shall be removed if there is an explicit
+   map clause.  */
+#define OMP_CLAUSE_MAP_IMPLICIT(NODE) \
+  (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)->base.default_def_flag)
 
 /* True on an OMP_CLAUSE_USE_DEVICE_PTR with an OpenACC 'if_present'
    clause.  */