[PATCH 11/14] OpenACC: Reimplement "inheritance" for lexically-nested offload regions

[Julian Brown <julian@codesourcery.com>]

This patch reimplements "lexical inheritance" for OpenACC offload regions
inside "data" regions, allowing e.g. this to work:

  int *ptr;
  [...]
  #pragma acc data copyin(ptr[10:2])
  {
    #pragma acc parallel
    { ... }
  }

here, the "copyin" is mirrored on the inner "acc parallel" as
"present(ptr[10:2])" -- allowing code within the parallel to use that
section of the array even though the mapping is implicit.

In terms of implementation, this works by expanding mapping nodes for
"acc data" to include pointer mappings that might be needed by inner
offload regions. The resulting mapping group is then copied to the inner
offload region as needed, rewriting the first node to "force_present".
The pointer mapping nodes are then removed from the "acc data" later
during gimplification.

For OpenMP, pointer mapping nodes on equivalent "omp data" regions are
not needed, so remain suppressed during expansion.

2023-06-16  Julian Brown  <julian@codesourcery.com>

gcc/c-family/
	* c-omp.cc (c_omp_address_inspector::expand_array_base): Don't omit
	pointer nodes for OpenACC.

gcc/
	* gimplify.cc (omp_tsort_mark, omp_mapping_group): Move before
	gimplify_omp_ctx. Add constructor to omp_mapping_group.
	(gimplify_omp_ctx): Add DECL_DATA_CLAUSE field.
	(new_omp_context, delete_omp_context): Initialise and free above field.
	(omp_gather_mapping_groups_1): Use constructor for omp_mapping_group.
	(gimplify_scan_omp_clauses): Record mappings that might be lexically
	inherited.  Don't remove
	GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE yet.
	(gomp_oacc_needs_data_present): New function.
	(gimplify_adjust_omp_clauses_1): Implement lexical inheritance
	behaviour for OpenACC.
	(gimplify_adjust_omp_clauses): Remove
	GOMP_MAP_FIRSTPRIVATE_POINTER/GOMP_MAP_FIRSTPRIVATE_REFERENCE here
	instead, after lexical inheritance is done.

gcc/testsuite/
	* c-c++-common/goacc/acc-data-chain.c: Re-enable scan test.
	* gfortran.dg/goacc/pr70828.f90: Likewise.
	* gfortran.dg/goacc/assumed-size.f90: New test.

libgomp/
	* testsuite/libgomp.oacc-c-c++-common/pr70828.c: Un-XFAIL.
	* testsuite/libgomp.oacc-c-c++-common/pr70828-2.c: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828.f90: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828-2.f90: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828-3.f90: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828-4.f90: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828-5.f90: Un-XFAIL.
	* testsuite/libgomp.oacc-fortran/pr70828-6.f90: Un-XFAIL.
---
 gcc/c-family/c-omp.cc                         |  13 +-
 gcc/gimplify.cc                               | 208 +++++++++++++-----
 .../c-c++-common/goacc/acc-data-chain.c       |   4 +-
 .../gfortran.dg/goacc/assumed-size.f90        |  35 +++
 gcc/testsuite/gfortran.dg/goacc/pr70828.f90   |   3 +-
 .../libgomp.oacc-c-c++-common/pr70828-2.c     |   2 -
 .../libgomp.oacc-c-c++-common/pr70828.c       |   2 -
 .../libgomp.oacc-fortran/pr70828-2.f90        |   2 -
 .../libgomp.oacc-fortran/pr70828-3.f90        |   2 -
 .../libgomp.oacc-fortran/pr70828-4.f90        |   2 -
 .../libgomp.oacc-fortran/pr70828-5.f90        |   2 -
 .../libgomp.oacc-fortran/pr70828-6.f90        |   2 -
 .../libgomp.oacc-fortran/pr70828.f90          |   2 -
 13 files changed, 202 insertions(+), 77 deletions(-)
 create mode 100644 gcc/testsuite/gfortran.dg/goacc/assumed-size.f90

diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index e55b2aec920..291a26293ef 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -4313,7 +4313,8 @@ c_omp_address_inspector::expand_array_base (tree c,
 	/* The code handling "firstprivatize_array_bases" in gimplify.cc is
 	   relevant here.  What do we need to create for arrays at this
 	   stage?  (This condition doesn't feel quite right.  FIXME?)  */
-	if (!target
+	if (openmp
+	    && !target
 	    && (TREE_CODE (TREE_TYPE (addr_tokens[i + 1]->expr))
 		== ARRAY_TYPE))
 	  break;
@@ -4324,7 +4325,7 @@ c_omp_address_inspector::expand_array_base (tree c,
 					   virtual_origin);
 	tree data_addr = omp_accessed_addr (addr_tokens, i + 1, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
-	if (decl_p && target)
+	if (decl_p && (!openmp || target))
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
 	else
 	  {
@@ -4375,9 +4376,11 @@ c_omp_address_inspector::expand_array_base (tree c,
 	tree data_addr = omp_accessed_addr (addr_tokens, last_access, expr);
 	c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
 	/* For OpenACC, use FIRSTPRIVATE_POINTER for decls even on non-compute
-	   regions (e.g. "acc data" constructs).  It'll be removed anyway in
-	   gimplify.cc, but doing it this way maintains diagnostic
-	   behaviour.  */
+	   regions (e.g. "acc data" constructs).  It is used during "lexical
+	   inheritance" of mapping clauses on enclosed target
+	   (parallel/serial/kernels) regions, i.e. creating "present" mappings
+	   for sections of pointer-based arrays.  It's also used for
+	   diagnostics.  */
 	if (decl_p && (target || !openmp) && !chain_p && !declare_target_p)
 	  OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
 	else
diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index e21e9d99cc9..55befa4d3c1 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -218,6 +218,48 @@ enum gimplify_defaultmap_kind
   GDMK_POINTER
 };
 
+/* Used for topological sorting of mapping groups.  UNVISITED means we haven't
+   started processing the group yet.  The TEMPORARY mark is used when we first
+   encounter a group on a depth-first traversal, and the PERMANENT mark is used
+   when we have processed all the group's children (i.e. all the base pointers
+   referred to by the group's mapping nodes, recursively).  */
+
+enum omp_tsort_mark {
+  UNVISITED,
+  TEMPORARY,
+  PERMANENT
+};
+
+/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map"
+   clause.  */
+
+struct omp_mapping_group {
+  tree *grp_start;
+  tree grp_end;
+  omp_tsort_mark mark;
+  /* If we've removed the group but need to reindex, mark the group as
+     deleted.  */
+  bool deleted;
+  /* The group points to an already-created "GOMP_MAP_STRUCT
+     GOMP_MAP_ATTACH_DETACH" pair.  */
+  bool reprocess_struct;
+  /* The group should use "zero-length" allocations for pointers that are not
+     mapped "to" on the same directive.  */
+  bool fragile;
+  struct omp_mapping_group *sibling;
+  struct omp_mapping_group *next;
+
+  omp_mapping_group (tree *_start, tree _end)
+    : grp_start (_start), grp_end (_end), mark (UNVISITED), deleted (false),
+      reprocess_struct (false), fragile (false), sibling (NULL), next (NULL)
+    {
+    }
+
+  omp_mapping_group ()
+    {
+    }
+};
+
 struct gimplify_omp_ctx
 {
   struct gimplify_omp_ctx *outer_context;
@@ -239,6 +281,7 @@ struct gimplify_omp_ctx
   bool in_for_exprs;
   bool ompacc;
   int defaultmap[5];
+  hash_map<tree, omp_mapping_group *> *decl_data_clause;
 };
 
 struct privatize_reduction
@@ -473,6 +516,7 @@ new_omp_context (enum omp_region_type region_type)
   c->defaultmap[GDMK_AGGREGATE] = GOVD_MAP;
   c->defaultmap[GDMK_ALLOCATABLE] = GOVD_MAP;
   c->defaultmap[GDMK_POINTER] = GOVD_MAP;
+  c->decl_data_clause = NULL;
 
   return c;
 }
@@ -485,6 +529,7 @@ delete_omp_context (struct gimplify_omp_ctx *c)
   splay_tree_delete (c->variables);
   delete c->privatized_types;
   c->loop_iter_var.release ();
+  delete c->decl_data_clause;
   XDELETE (c);
 }
 
@@ -8988,18 +9033,6 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp,
   return base;
 }
 
-/* Used for topological sorting of mapping groups.  UNVISITED means we haven't
-   started processing the group yet.  The TEMPORARY mark is used when we first
-   encounter a group on a depth-first traversal, and the PERMANENT mark is used
-   when we have processed all the group's children (i.e. all the base pointers
-   referred to by the group's mapping nodes, recursively).  */
-
-enum omp_tsort_mark {
-  UNVISITED,
-  TEMPORARY,
-  PERMANENT
-};
-
 /* Hash for trees based on operand_equal_p.  Like tree_operand_hash
    but ignores side effects in the equality comparisons.  */
 
@@ -9016,26 +9049,6 @@ tree_operand_hash_no_se::equal (const value_type &t1,
   return operand_equal_p (t1, t2, OEP_MATCH_SIDE_EFFECTS);
 }
 
-/* A group of OMP_CLAUSE_MAP nodes that correspond to a single "map"
-   clause.  */
-
-struct omp_mapping_group {
-  tree *grp_start;
-  tree grp_end;
-  omp_tsort_mark mark;
-  /* If we've removed the group but need to reindex, mark the group as
-     deleted.  */
-  bool deleted;
-  /* The group points to an already-created "GOMP_MAP_STRUCT
-     GOMP_MAP_ATTACH_DETACH" pair.  */
-  bool reprocess_struct;
-  /* The group should use "zero-length" allocations for pointers that are not
-     mapped "to" on the same directive.  */
-  bool fragile;
-  struct omp_mapping_group *sibling;
-  struct omp_mapping_group *next;
-};
-
 DEBUG_FUNCTION void
 debug_mapping_group (omp_mapping_group *grp)
 {
@@ -9276,16 +9289,7 @@ omp_gather_mapping_groups_1 (tree *list_p, vec<omp_mapping_group> *groups,
 	continue;
 
       tree *grp_last_p = omp_group_last (cp);
-      omp_mapping_group grp;
-
-      grp.grp_start = cp;
-      grp.grp_end = *grp_last_p;
-      grp.mark = UNVISITED;
-      grp.sibling = NULL;
-      grp.deleted = false;
-      grp.reprocess_struct = false;
-      grp.fragile = false;
-      grp.next = NULL;
+      omp_mapping_group grp (cp, *grp_last_p);
       groups->safe_push (grp);
 
       cp = grp_last_p;
@@ -12267,6 +12271,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 		}
 	      break;
 	    }
+	  if (code == OACC_DATA && *grp_start_p != grp_end)
+	    {
+	      if (!ctx->decl_data_clause)
+		ctx->decl_data_clause = new hash_map<tree, omp_mapping_group *>;
+
+	      omp_mapping_group *grp
+		= new omp_mapping_group (grp_start_p, grp_end);
+
+	      gcc_assert (DECL_P (decl));
+
+	      ctx->decl_data_clause->put (decl, grp);
+	    }
 	  flags = GOVD_MAP | GOVD_EXPLICIT;
 	  if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
 	      || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
@@ -12953,11 +12969,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	  gcc_unreachable ();
 	}
 
-      if (code == OACC_DATA
-	  && 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))
-	remove = true;
       if (remove)
 	*list_p = OMP_CLAUSE_CHAIN (c);
       else
@@ -13098,6 +13109,52 @@ struct gimplify_adjust_omp_clauses_data
   gimple_seq *pre_p;
 };
 
+/* For OpenACC offload regions, the implicit data mappings for arrays must
+   respect explicit data clauses set by a containing acc data region.
+   Specifically, an array section on the data clause must be transformed into
+   an equivalent PRESENT mapping on the inner offload region.
+   This function returns a pointer to a mapping group if an array slice of DECL
+   is specified on a lexically-enclosing data construct, or returns NULL
+   otherwise.  */
+
+static omp_mapping_group *
+gomp_oacc_needs_data_present (tree decl)
+{
+  gimplify_omp_ctx *ctx = NULL;
+
+  if (gimplify_omp_ctxp->region_type != ORT_ACC_PARALLEL
+      && gimplify_omp_ctxp->region_type != ORT_ACC_SERIAL
+      && gimplify_omp_ctxp->region_type != ORT_ACC_KERNELS)
+    return NULL;
+
+  if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE
+      && TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+      && TREE_CODE (TREE_TYPE (decl)) != RECORD_TYPE
+      && (TREE_CODE (TREE_TYPE (decl)) != POINTER_TYPE
+	  || TREE_CODE (TREE_TYPE (TREE_TYPE (decl))) != ARRAY_TYPE))
+    return NULL;
+
+  decl = get_base_address (decl);
+
+  for (ctx = gimplify_omp_ctxp->outer_context; ctx; ctx = ctx->outer_context)
+    {
+      splay_tree_node on
+	= splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
+
+      if (ctx->region_type == ORT_ACC_DATA
+	  && on
+	  && (((int) on->value) & GOVD_EXPLICIT)
+	  && ctx->decl_data_clause != NULL)
+	{
+	  omp_mapping_group **pgrp = ctx->decl_data_clause->get (decl);
+	  if (pgrp)
+	    return *pgrp;
+	}
+    }
+
+  return NULL;
+}
+
 /* For all variables that were not actually used within the context,
    remove PRIVATE, SHARED, and FIRSTPRIVATE clauses.  */
 
@@ -13219,6 +13276,7 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
   clause = build_omp_clause (input_location, code);
   OMP_CLAUSE_DECL (clause) = decl;
   OMP_CLAUSE_CHAIN (clause) = chain;
+  omp_mapping_group *outer_grp;
   if (private_debug)
     OMP_CLAUSE_PRIVATE_DEBUG (clause) = 1;
   else if (code == OMP_CLAUSE_PRIVATE && (flags & GOVD_PRIVATE_OUTER_REF))
@@ -13227,6 +13285,58 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
 	   && (flags & GOVD_WRITTEN) == 0
 	   && omp_shared_to_firstprivate_optimizable_decl_p (decl))
     OMP_CLAUSE_SHARED_READONLY (clause) = 1;
+  else if ((gimplify_omp_ctxp->region_type & ORT_ACC) != 0
+	   && (code == OMP_CLAUSE_MAP || code == OMP_CLAUSE_FIRSTPRIVATE)
+	   && (outer_grp = gomp_oacc_needs_data_present (decl)))
+    {
+      if (code == OMP_CLAUSE_FIRSTPRIVATE)
+	/* Oops, we have the wrong type of clause.  Rebuild it.  */
+	clause = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+				   OMP_CLAUSE_MAP);
+
+      tree mapping = *outer_grp->grp_start;
+
+      OMP_CLAUSE_SET_MAP_KIND (clause, GOMP_MAP_FORCE_PRESENT);
+      OMP_CLAUSE_DECL (clause) = unshare_expr (OMP_CLAUSE_DECL (mapping));
+      OMP_CLAUSE_SIZE (clause) = unshare_expr (OMP_CLAUSE_SIZE (mapping));
+
+      /* Copy subsequent nodes (that are part of the mapping group) after the
+	 initial one from the outer "acc data" directive -- "pointer" nodes,
+	 including firstprivate_reference, pointer sets, etc.  */
+
+      tree ptr = OMP_CLAUSE_CHAIN (mapping);
+      tree *ins = &OMP_CLAUSE_CHAIN (clause);
+      tree sentinel = OMP_CLAUSE_CHAIN (outer_grp->grp_end);
+      for (; ptr && ptr != sentinel; ptr = OMP_CLAUSE_CHAIN (ptr))
+	{
+	  tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+				      OMP_CLAUSE_MAP);
+	  OMP_CLAUSE_SET_MAP_KIND (nc, OMP_CLAUSE_MAP_KIND (ptr));
+	  OMP_CLAUSE_DECL (nc) = unshare_expr (OMP_CLAUSE_DECL (ptr));
+	  OMP_CLAUSE_SIZE (nc) = unshare_expr (OMP_CLAUSE_SIZE (ptr));
+	  *ins = nc;
+	  ins = &OMP_CLAUSE_CHAIN (nc);
+	}
+
+      *ins = chain;
+
+      gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+      gimplify_omp_ctxp = ctx->outer_context;
+      for (ptr = clause; ptr != chain; ptr = OMP_CLAUSE_CHAIN (ptr))
+	{
+	  /* The condition is specifically to not gimplify here if we have a
+	     DECL_P with a DECL_VALUE_EXPR -- i.e. a VLA, or variable-sized
+	     array section.  If we do, omp-low.cc does not see the DECL_P it
+	     expects here for e.g. firstprivate_pointer or
+	     firstprivate_reference.  */
+	  if (!DECL_P (OMP_CLAUSE_DECL (ptr)))
+	    gimplify_expr (&OMP_CLAUSE_DECL (ptr), pre_p, NULL,
+			   is_gimple_lvalue, fb_lvalue);
+	  gimplify_expr (&OMP_CLAUSE_SIZE (ptr), pre_p, NULL,
+			 is_gimple_val, fb_rvalue);
+	}
+      gimplify_omp_ctxp = ctx;
+    }
   else if (code == OMP_CLAUSE_FIRSTPRIVATE && (flags & GOVD_EXPLICIT) == 0)
     OMP_CLAUSE_FIRSTPRIVATE_IMPLICIT (clause) = 1;
   else if (code == OMP_CLAUSE_MAP && (flags & GOVD_MAP_0LEN_ARRAY) != 0)
@@ -13689,16 +13799,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	    case OMP_TARGET:
 	      break;
 	    case OACC_DATA:
-	      if (TREE_CODE (TREE_TYPE (decl)) != ARRAY_TYPE)
-		break;
-	      goto check_firstprivate;
 	    case OACC_ENTER_DATA:
 	    case OACC_EXIT_DATA:
 	    case OMP_TARGET_DATA:
 	    case OMP_TARGET_ENTER_DATA:
 	    case OMP_TARGET_EXIT_DATA:
 	    case OACC_HOST_DATA:
-	    check_firstprivate:
 	      if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
 		  || (OMP_CLAUSE_MAP_KIND (c)
 		      == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
diff --git a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
index 932786cec76..622f1992f88 100644
--- a/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
+++ b/gcc/testsuite/c-c++-common/goacc/acc-data-chain.c
@@ -21,6 +21,4 @@ int main(int argc, char *argv[])
 }
 
 // { dg-final { scan-tree-dump-times "omp target oacc_data map\\(from:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(to:a\\\[0\\\] \\\[len: 400\\\]\\)" 1 "gimple" } }
-/* This isn't expected to work while the "lexical inheritance" support is
-   reverted.  */
-// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map.alloc:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(alloc:a \\\[pointer assign, bias: 0\\\]\\)" 0 "gimple" } }
+// { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:b\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:b \\\[pointer assign, bias: 0\\\]\\) map\\(force_present:a\\\[0\\\] \\\[len: 400\\\]\\) map\\(firstprivate:a \\\[pointer assign, bias: 0\\\]\\)" 1 "gimple" } }
diff --git a/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90 b/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90
new file mode 100644
index 00000000000..4fced2e70c9
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/assumed-size.f90
@@ -0,0 +1,35 @@
+! Test if implicitly determined data clauses work with an
+! assumed-sized array variable.  Note that the array variable, 'a',
+! has been explicitly copied in and out via acc enter data and acc
+! exit data, respectively.
+
+! This does not appear to be supported by the OpenACC standard as of version
+! 3.0.  Check for an appropriate error message.
+
+program test
+  implicit none
+
+  integer, parameter :: n = 100
+  integer a(n), i
+
+  call dtest (a, n)
+
+  do i = 1, n
+     if (a(i) /= i) call abort
+  end do
+end program test
+
+subroutine dtest (a, n)
+  integer i, n
+  integer a(*)
+
+  !$acc enter data copyin(a(1:n))
+
+  !$acc parallel loop
+! { dg-error {implicit mapping of assumed size array 'a'} "" { target *-*-* } .-1 }
+  do i = 1, n
+     a(i) = i
+  end do
+
+  !$acc exit data copyout(a(1:n))
+end subroutine dtest
diff --git a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90 b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90
index 72b0d9ae92c..fcfe0865fc4 100644
--- a/gcc/testsuite/gfortran.dg/goacc/pr70828.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/pr70828.f90
@@ -19,5 +19,4 @@ program test
 end program test
 
 ! { dg-final { scan-tree-dump-times "omp target oacc_data map\\(tofrom:data\\\[\_\[0-9\]+\\\] \\\[len: _\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: _\[0-9\]+\\\]\\)" 1 "gimple" } }
-! Disable for now
-! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 0 "gimple" } }
+! { dg-final { scan-tree-dump-times "omp target oacc_parallel map\\(force_present:data\\\[D\\.\[0-9\]+\\\] \\\[len: D\\.\[0-9\]+\\\]\\) map\\(alloc:data \\\[pointer assign, bias: D\\.\[0-9\]+\\\]\\)" 1 "gimple" } }
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c
index da5bb3f93c3..357114ccfd3 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828-2.c
@@ -32,5 +32,3 @@ main (int argc, char* argv[])
 
   return 0;
 }
-
-/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
index 85d09bff1df..4b6dbd7538f 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/pr70828.c
@@ -25,5 +25,3 @@ main ()
 
   return 0;
 }
-
-/* { dg-xfail-run-if "PR70828" { ! openacc_host_selected } } */
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
index 2892b3d5938..22a956622bb 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-2.f90
@@ -29,5 +29,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
index e28193b1a22..ff17d10cfa3 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-3.f90
@@ -32,5 +32,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
index 918295d5c8b..01da999b33d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-4.f90
@@ -29,5 +29,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
index 3b5d05d1379..8a16e3d5550 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-5.f90
@@ -27,5 +27,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
index d48168b22eb..e99c3649159 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828-6.f90
@@ -26,5 +26,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90 b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
index 5db49e1a569..f87d232fe42 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/pr70828.f90
@@ -22,5 +22,3 @@ program test
      end if
   end do
 end program test
-
-! { dg-xfail-run-if "PR70828" { ! openacc_host_selected } }
-- 
2.31.1