[gcc r13-2717] OpenMP: Push attaches to end of clause list in "target" regions

[gcc r13-2717] OpenMP: Push attaches to end of clause list in "target" regions

[Julian Brown]

https://gcc.gnu.org/g:313879d8768d08dea035efd7fd62b753dc91c364

commit r13-2717-g313879d8768d08dea035efd7fd62b753dc91c364
Author: Julian Brown <julian@codesourcery.com>
Date:   Wed Sep 7 21:45:07 2022 +0000

    OpenMP: Push attaches to end of clause list in "target" regions
    
    This patch moves GOMP_MAP_ATTACH{_ZERO_LENGTH_ARRAY_SECTION} nodes to
    the end of the clause list, for offload regions.  This ensures that when
    we do the attach operation, both the "attachment point" and the target
    region have both already been mapped on the target.  This avoids a
    pathological case that can otherwise happen with struct sibling-list
    handling.
    
    This version of the patch moves the attach-node motion to
    gimplify_adjust_omp_clauses.
    
    2022-09-15  Julian Brown  <julian@codesourcery.com>
    
    gcc/
            * gimplify.cc (omp_segregate_mapping_groups): Update comment.
            (gimplify_adjust_omp_clauses): Move ATTACH and
            ATTACH_ZERO_LENGTH_ARRAY_SECTION nodes to the end of the clause list
            for offloaded OpenMP regions.
    
    gcc/testsuite/
            * g++.dg/gomp/target-lambda-1.C: Adjust expected scan output.

Diff:
---
 gcc/gimplify.cc                             | 37 ++++++++++++++++++++++++++++-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C |  2 +-
 2 files changed, 37 insertions(+), 2 deletions(-)

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index 2ae0c8cb250..4d032c6bf06 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9639,7 +9639,9 @@ omp_tsort_mapping_groups (vec<omp_mapping_group> *groups,
 /* Split INLIST into two parts, moving groups corresponding to
    ALLOC/RELEASE/DELETE mappings to one list, and other mappings to another.
    The former list is then appended to the latter.  Each sub-list retains the
-   order of the original list.  */
+   order of the original list.
+   Note that ATTACH nodes are later moved to the end of the list in
+   gimplify_adjust_omp_clauses, for target regions.  */
 
 static omp_mapping_group *
 omp_segregate_mapping_groups (omp_mapping_group *inlist)
@@ -12411,10 +12413,15 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	    *list_p = c2;
 	  }
     }
+
+  tree attach_list = NULL_TREE;
+  tree *attach_tail = &attach_list;
+
   while ((c = *list_p) != NULL)
     {
       splay_tree_node n;
       bool remove = false;
+      bool move_attach = false;
 
       switch (OMP_CLAUSE_CODE (c))
 	{
@@ -12576,6 +12583,19 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 	      remove = true;
 	      break;
 	    }
+	  /* If we have a target region, we can push all the attaches to the
+	     end of the list (we may have standalone "attach" operations
+	     synthesized for GOMP_MAP_STRUCT nodes that must be processed after
+	     the attachment point AND the pointed-to block have been mapped).
+	     If we have something else, e.g. "enter data", we need to keep
+	     "attach" nodes together with the previous node they attach to so
+	     that separate "exit data" operations work properly (see
+	     libgomp/target.c).  */
+	  if ((ctx->region_type & ORT_TARGET) != 0
+	      && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+		  || (OMP_CLAUSE_MAP_KIND (c)
+		      == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)))
+	    move_attach = true;
 	  decl = OMP_CLAUSE_DECL (c);
 	  /* Data clauses associated with reductions must be
 	     compatible with present_or_copy.  Warn and adjust the clause
@@ -12890,10 +12910,25 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
 
       if (remove)
 	*list_p = OMP_CLAUSE_CHAIN (c);
+      else if (move_attach)
+	{
+	  /* Remove attach node from here, separate out into its own list.  */
+	  *attach_tail = c;
+	  *list_p = OMP_CLAUSE_CHAIN (c);
+	  OMP_CLAUSE_CHAIN (c) = NULL_TREE;
+	  attach_tail = &OMP_CLAUSE_CHAIN (c);
+	}
       else
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
+  /* Splice attach nodes at the end of the list.  */
+  if (attach_list)
+    {
+      *list_p = attach_list;
+      list_p = attach_tail;
+    }
+
   /* Add in any implicit data sharing.  */
   struct gimplify_adjust_omp_clauses_data data;
   if ((gimplify_omp_ctxp->region_type & ORT_ACC) == 0)
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
index bff7fa7c669..5ce8ceadb19 100644
--- a/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-1.C
@@ -87,7 +87,7 @@ int main (void)
   return 0;
 }
 
-/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(b\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:iptr \[pointer assign, bias: 0\]\) map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\) map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:__closure->__iptr \[bias: 0\]\) map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)} "gimple" } } */
 
 /* { dg-final { scan-tree-dump {#pragma omp target num_teams.* firstprivate\(end\) firstprivate\(begin\) map\(to:loop \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: 0\]\) map\(attach_zero_length_array_section:loop\.__data1 \[bias: 0\]\)} "gimple" } } */