public inbox for fortran@gcc.gnu.org
 help / color / mirror / Atom feed
From: Julian Brown <julian@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: <fortran@gcc.gnu.org>, Jakub Jelinek <jakub@redhat.com>,
	<tobias@codesourcery.com>, <cltang@codesourcery.com>
Subject: [PATCH v3 05/11] OpenMP: push attaches to end of clause list in "target" regions
Date: Tue, 13 Sep 2022 14:03:15 -0700	[thread overview]
Message-ID: <479bff9d51ee4db1ff46e0edaaf24d2a601f7a0d.1663101299.git.julian@codesourcery.com> (raw)
In-Reply-To: <cover.1663101299.git.julian@codesourcery.com>

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.

2022-09-13  Julian Brown  <julian@codesourcery.com>

gcc/
	* gimplify.cc (omp_segregate_mapping_groups): Update comment.
	(omp_push_attaches_to_end): New function.
	(gimplify_scan_omp_clauses): Use omp_push_attaches_to_end for offloaded
	regions.
---
 gcc/gimplify.cc                             | 66 ++++++++++++++++++++-
 gcc/testsuite/g++.dg/gomp/target-lambda-1.C |  2 +-
 2 files changed, 65 insertions(+), 3 deletions(-)

diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc
index c7998c2ccbd..bc7848843b3 100644
--- a/gcc/gimplify.cc
+++ b/gcc/gimplify.cc
@@ -9641,7 +9641,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.
+   See also omp_push_attaches_to_end below -- we call that later after scanning
+   omp clauses.  */
 
 static omp_mapping_group *
 omp_segregate_mapping_groups (omp_mapping_group *inlist)
@@ -9681,6 +9683,55 @@ omp_segregate_mapping_groups (omp_mapping_group *inlist)
   return tf_groups;
 }
 
+/* This function 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, 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.
+
+   Do not call this for non-offload regions, e.g. for "enter data" or
+   "exit data" directives.
+
+   The order of attach nodes and of non-attach nodes is otherwise retained.  */
+
+static tree
+omp_push_attaches_to_end (tree list)
+{
+  tree nonattach_list = NULL_TREE, attach_list = NULL_TREE;
+  tree *nonattach_tail = &nonattach_list, *attach_tail = &attach_list;
+
+  for (tree w = list; w;)
+    {
+      tree next = OMP_CLAUSE_CHAIN (w);
+
+      if (OMP_CLAUSE_CODE (w) != OMP_CLAUSE_MAP)
+	goto nonattach;
+
+      switch (OMP_CLAUSE_MAP_KIND (w))
+	{
+	case GOMP_MAP_ATTACH:
+	case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+	  *attach_tail = w;
+	  OMP_CLAUSE_CHAIN (w) = NULL_TREE;
+	  attach_tail = &OMP_CLAUSE_CHAIN (w);
+	  break;
+
+	default:
+	nonattach:
+	  *nonattach_tail = w;
+	  OMP_CLAUSE_CHAIN (w) = NULL_TREE;
+	  nonattach_tail = &OMP_CLAUSE_CHAIN (w);
+	}
+
+      w = next;
+    }
+
+  /* Splice lists together.  */
+  *nonattach_tail = attach_list;
+
+  return nonattach_list;
+}
+
 /* Given a list LIST_P containing groups of mappings given by GROUPS, reorder
    those groups based on the output list of omp_tsort_mapping_groups --
    singly-linked, threaded through each element's NEXT pointer starting at
@@ -11950,7 +12001,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
 	list_p = &OMP_CLAUSE_CHAIN (c);
     }
 
-  ctx->clauses = *orig_list_p;
+  if ((region_type & ORT_TARGET) != 0)
+    /* 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).  */
+    ctx->clauses = omp_push_attaches_to_end (*orig_list_p);
+  else
+    /* ...but 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).  */
+    ctx->clauses = *orig_list_p;
+
   gimplify_omp_ctxp = ctx;
 }
 
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" } } */
 
-- 
2.29.2


  parent reply	other threads:[~2022-09-13 21:03 UTC|newest]

Thread overview: 36+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2022-09-13 21:01 [PATCH v3 00/11] OpenMP 5.0: Struct & mapping clause expansion rework Julian Brown
2022-09-13 21:01 ` [PATCH v3 01/11] OpenMP 5.0: Clause ordering for OpenMP 5.0 (topological sorting by base pointer) Julian Brown
2022-09-14 10:34   ` Jakub Jelinek
2022-09-13 21:01 ` [PATCH v3 02/11] Remove omp_target_reorder_clauses Julian Brown
2022-09-14 10:35   ` Jakub Jelinek
2022-09-13 21:01 ` [PATCH v3 03/11] OpenMP/OpenACC struct sibling list gimplification extension and rework Julian Brown
2022-09-14 11:21   ` Jakub Jelinek
2022-09-13 21:01 ` [PATCH v3 04/11] OpenMP/OpenACC: mapping group list-handling improvements Julian Brown
2022-09-14 11:30   ` Jakub Jelinek
2022-09-13 21:03 ` Julian Brown [this message]
2022-09-14 12:44   ` [PATCH v3 05/11] OpenMP: push attaches to end of clause list in "target" regions Jakub Jelinek
2022-09-18 19:10     ` Julian Brown
2022-09-18 19:18       ` Jakub Jelinek
2022-09-13 21:03 ` [PATCH v3 06/11] OpenMP: Pointers and member mappings Julian Brown
2022-09-14 12:53   ` Jakub Jelinek
2022-09-18 19:19     ` Julian Brown
2022-09-22 13:17       ` Jakub Jelinek
2022-09-23  7:29         ` Julian Brown
2022-09-23  9:38           ` Jakub Jelinek
2022-09-23 12:10           ` Tobias Burnus
2022-09-30 13:30             ` Julian Brown
2022-09-30 14:42               ` Tobias Burnus
2022-09-30 15:01               ` Tobias Burnus
2022-09-13 21:03 ` [PATCH v3 07/11] OpenMP/OpenACC: Reindent TO/FROM/_CACHE_ stanza in {c_}finish_omp_clause Julian Brown
2022-09-14 13:06   ` Jakub Jelinek
2022-09-13 21:03 ` [PATCH v3 08/11] OpenMP/OpenACC: Rework clause expansion and nested struct handling Julian Brown
2022-09-14 13:24   ` Jakub Jelinek
2022-09-14 13:59     ` Julian Brown
2022-09-19 19:40     ` Julian Brown
2022-09-22 13:20       ` Jakub Jelinek
2022-09-13 21:03 ` [PATCH v3 09/11] FYI/unfinished: OpenMP: lvalue parsing for map clauses (C++) Julian Brown
2022-09-13 21:04 ` [PATCH v3 10/11] Use OMP_ARRAY_SECTION instead of TREE_LIST in C++ FE Julian Brown
2022-09-13 21:04 ` [PATCH v3 11/11] FYI/unfinished: OpenMP 5.0 "declare mapper" support for C++ Julian Brown
2022-09-14  6:30   ` FYI: "declare mapper" patch set for Fortran (June 2022) (was: [PATCH v3 11/11] FYI/unfinished: OpenMP 5.0 "declare mapper" support for C++) Tobias Burnus
2022-09-14 14:58   ` [PATCH v3 11/11] FYI/unfinished: OpenMP 5.0 "declare mapper" support for C++ Jakub Jelinek
2022-09-14 16:32     ` Julian Brown

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=479bff9d51ee4db1ff46e0edaaf24d2a601f7a0d.1663101299.git.julian@codesourcery.com \
    --to=julian@codesourcery.com \
    --cc=cltang@codesourcery.com \
    --cc=fortran@gcc.gnu.org \
    --cc=gcc-patches@gcc.gnu.org \
    --cc=jakub@redhat.com \
    --cc=tobias@codesourcery.com \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).