Hi! On 2019-12-17T22:03:48-0800, Julian Brown <jul...@codesourcery.com> wrote: > This part contains the middle-end support for OpenACC 2.6 attach and > detach operations, either as standalone clauses or as "attach/detach" > actions triggered by other (data movement) clauses, as detailed in the > specification.
As mentioned in <https://gcc.gnu.org/PR93026>, "that commit [r279626] is doing more than just the OpenACC 2.6 manual deep copy changes; see the 'gcc/gimplify.c' changes related to the PR92929 discussion" etc. (see <https://gcc.gnu.org/PR92929>). See attached "[PR93026, PR92929] Adjust 'gfortran.dg/goacc/finalize-1.f' for r279626 changes"; committed to trunk in r279700. Grüße Thomas > ChangeLog > > gcc/ > * gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS. > (insert_struct_comp_map): Support derived-type member mappings > for arrays with descriptors which use GOMP_MAP_TO_PSET. Support > GOMP_MAP_ATTACH_DETACH. > (gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA > mappings. Handle attach/detach clauses and component references. > (gimplify_adjust_omp_clauses_1): Skip adjustments for explicit > attach/detach clauses. > (gimplify_omp_target_update): Handle finalize for detach. > * omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH, > GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH. > * tree-pretty-print.c (dump_omp_clause): Likewise, plus > GOMP_MAP_ATTACH_DETACH. > > include/ > * gomp-constants.h (gomp_map_kind): Add GOMP_MAP_ATTACH_DETACH. > --- > gcc/gimplify.c | 232 ++++++++++++++++++++++++++++++++++----- > gcc/omp-low.c | 3 + > gcc/tree-pretty-print.c | 18 +++ > include/gomp-constants.h | 6 +- > 4 files changed, 229 insertions(+), 30 deletions(-) > > diff --git a/gcc/gimplify.c b/gcc/gimplify.c > index e3088dcbe05..e3d5bc83c4f 100644 > --- a/gcc/gimplify.c > +++ b/gcc/gimplify.c > @@ -123,6 +123,10 @@ enum gimplify_omp_var_data > /* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause. */ > GOVD_REDUCTION_INSCAN = 0x2000000, > > + /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for > + fields. */ > + GOVD_MAP_HAS_ATTACHMENTS = 8388608, > + > GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE > | GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR > | GOVD_LOCAL) > @@ -8209,20 +8213,33 @@ insert_struct_comp_map (enum tree_code code, tree c, > tree struct_node, > tree prev_node, tree *scp) > { > enum gomp_map_kind mkind > - = code == OMP_TARGET_EXIT_DATA ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; > + = (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA) > + ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; > > tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); > tree cl = scp ? prev_node : c2; > OMP_CLAUSE_SET_MAP_KIND (c2, mkind); > OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c)); > OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node; > - OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node); > + if (OMP_CLAUSE_CHAIN (prev_node) != c > + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP > + && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) > + == GOMP_MAP_TO_PSET)) > + OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node)); > + else > + OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node); > if (struct_node) > OMP_CLAUSE_CHAIN (struct_node) = c2; > > /* We might need to create an additional mapping if we have a reference to > a > - pointer (in C++). */ > - if (OMP_CLAUSE_CHAIN (prev_node) != c) > + pointer (in C++). Don't do this if we have something other than a > + GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET. */ > + if (OMP_CLAUSE_CHAIN (prev_node) != c > + && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP > + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) > + == GOMP_MAP_ALWAYS_POINTER) > + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node)) > + == GOMP_MAP_ATTACH_DETACH))) > { > tree c4 = OMP_CLAUSE_CHAIN (prev_node); > tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP); > @@ -8329,6 +8346,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > struct gimplify_omp_ctx *ctx, *outer_ctx; > tree c; > hash_map<tree, tree> *struct_map_to_clause = NULL; > + hash_set<tree> *struct_deref_set = NULL; > tree *prev_list_p = NULL, *orig_list_p = list_p; > int handled_depend_iterators = -1; > int nowait = -1; > @@ -8731,8 +8749,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > case OMP_TARGET_DATA: > case OMP_TARGET_ENTER_DATA: > case OMP_TARGET_EXIT_DATA: > - case OACC_ENTER_DATA: > - case OACC_EXIT_DATA: > case OACC_HOST_DATA: > if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER > || (OMP_CLAUSE_MAP_KIND (c) > @@ -8741,6 +8757,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > mapped, but not the pointer to it. */ > remove = true; > break; > + case OACC_ENTER_DATA: > + case OACC_EXIT_DATA: > + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER > + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET > + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER > + || (OMP_CLAUSE_MAP_KIND (c) > + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) > + remove = true; > + break; > default: > break; > } > @@ -8814,7 +8839,35 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > pd = &TREE_OPERAND (decl, 0); > decl = TREE_OPERAND (decl, 0); > } > - if (TREE_CODE (decl) == COMPONENT_REF) > + bool indir_p = false; > + tree orig_decl = decl; > + tree decl_ref = NULL_TREE; > + if ((region_type & ORT_ACC) != 0 > + && TREE_CODE (*pd) == COMPONENT_REF > + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH > + && code != OACC_UPDATE) > + { > + 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)) > + { > + indir_p = true; > + decl = TREE_OPERAND (decl, 0); > + } > + if (TREE_CODE (decl) == INDIRECT_REF > + && DECL_P (TREE_OPERAND (decl, 0)) > + && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0))) > + == REFERENCE_TYPE)) > + { > + decl_ref = decl; > + decl = TREE_OPERAND (decl, 0); > + } > + } > + } > + else if (TREE_CODE (decl) == COMPONENT_REF) > { > while (TREE_CODE (decl) == COMPONENT_REF) > decl = TREE_OPERAND (decl, 0); > @@ -8824,13 +8877,76 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > == REFERENCE_TYPE)) > decl = TREE_OPERAND (decl, 0); > } > + if (decl != orig_decl && DECL_P (decl) && indir_p) > + { > + gomp_map_kind k = (code == OACC_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. */ > + OMP_CLAUSE_SET_MAP_KIND (c, k); > + flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT; > + tree next_clause = OMP_CLAUSE_CHAIN (c); > + if (k == GOMP_MAP_ATTACH > + && code != OACC_ENTER_DATA > + && (!next_clause > + || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP) > + || (OMP_CLAUSE_MAP_KIND (next_clause) > + != GOMP_MAP_POINTER) > + || OMP_CLAUSE_DECL (next_clause) != decl) > + && (!struct_deref_set > + || !struct_deref_set->contains (decl))) > + { > + if (!struct_deref_set) > + struct_deref_set = new hash_set<tree> (); > + /* As well as the attach, we also need a > + FIRSTPRIVATE_POINTER clause to properly map the > + pointer to the struct base. */ > + tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), > + OMP_CLAUSE_MAP); > + OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC); > + OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2) > + = 1; > + tree charptr_zero > + = build_int_cst (build_pointer_type (char_type_node), > + 0); > + OMP_CLAUSE_DECL (c2) > + = build2 (MEM_REF, char_type_node, > + decl_ref ? decl_ref : decl, charptr_zero); > + OMP_CLAUSE_SIZE (c2) = size_zero_node; > + tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), > + OMP_CLAUSE_MAP); > + OMP_CLAUSE_SET_MAP_KIND (c3, > + GOMP_MAP_FIRSTPRIVATE_POINTER); > + OMP_CLAUSE_DECL (c3) = decl; > + OMP_CLAUSE_SIZE (c3) = size_zero_node; > + tree mapgrp = *prev_list_p; > + *prev_list_p = c2; > + OMP_CLAUSE_CHAIN (c3) = mapgrp; > + OMP_CLAUSE_CHAIN (c2) = c3; > + > + struct_deref_set->add (decl); > + } > + goto do_add_decl; > + } > + /* An "attach/detach" operation on an update directive should > + behave as a GOMP_MAP_ALWAYS_POINTER. Beware that > + unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER > + depends on the previous mapping. */ > + 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)) > + 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) > { > if (error_operand_p (decl)) > { > @@ -8851,7 +8967,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > break; > } > > - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER) > + if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER > + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH) > { > /* Error recovery. */ > if (prev_list_p == NULL) > @@ -8884,20 +9001,47 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > = splay_tree_lookup (ctx->variables, (splay_tree_key)decl); > bool ptr = (OMP_CLAUSE_MAP_KIND (c) > == GOMP_MAP_ALWAYS_POINTER); > + bool attach_detach = (OMP_CLAUSE_MAP_KIND (c) > + == GOMP_MAP_ATTACH_DETACH); > + bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH > + || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH; > + bool has_attachments = false; > + /* For OpenACC, pointers in structs should trigger an > + attach action. */ > + if (attach_detach && (region_type & ORT_ACC) != 0) > + { > + /* 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; > + OMP_CLAUSE_SET_MAP_KIND (c, k); > + has_attachments = true; > + } > if (n == NULL || (n->value & GOVD_MAP) == 0) > { > tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c), > OMP_CLAUSE_MAP); > - OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT); > + gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT > + : GOMP_MAP_STRUCT; > + > + OMP_CLAUSE_SET_MAP_KIND (l, k); > if (base_ref) > OMP_CLAUSE_DECL (l) = unshare_expr (base_ref); > else > OMP_CLAUSE_DECL (l) = decl; > - OMP_CLAUSE_SIZE (l) = size_int (1); > + OMP_CLAUSE_SIZE (l) > + = (!attach > + ? size_int (1) > + : DECL_P (OMP_CLAUSE_DECL (l)) > + ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l)) > + : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l)))); > if (struct_map_to_clause == NULL) > struct_map_to_clause = new hash_map<tree, tree>; > struct_map_to_clause->put (decl, l); > - if (ptr) > + if (ptr || attach_detach) > { > insert_struct_comp_map (code, c, l, *prev_list_p, > NULL); > @@ -8923,23 +9067,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > OMP_CLAUSE_CHAIN (l) = c2; > } > flags = GOVD_MAP | GOVD_EXPLICIT; > - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) > + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) > + || ptr > + || attach_detach) > flags |= GOVD_SEEN; > + if (has_attachments) > + flags |= GOVD_MAP_HAS_ATTACHMENTS; > goto do_add_decl; > } > - else > + else if (struct_map_to_clause) > { > tree *osc = struct_map_to_clause->get (decl); > tree *sc = NULL, *scp = NULL; > - if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr) > + if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) > + || ptr > + || attach_detach) > n->value |= GOVD_SEEN; > sc = &OMP_CLAUSE_CHAIN (*osc); > if (*sc != c > && (OMP_CLAUSE_MAP_KIND (*sc) > - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) > + == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) > sc = &OMP_CLAUSE_CHAIN (*sc); > + /* Here "prev_list_p" is the end of the inserted > + alloc/release nodes after the struct node, OSC. */ > for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc)) > - if (ptr && sc == prev_list_p) > + if ((ptr || attach_detach) && sc == prev_list_p) > break; > else if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) > != COMPONENT_REF > @@ -8992,7 +9144,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > || (known_eq (offset1, offsetn) > && maybe_lt (bitpos1, bitposn))) > { > - if (ptr) > + if (ptr || attach_detach) > scp = sc; > else > break; > @@ -9000,10 +9152,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > } > if (remove) > break; > - OMP_CLAUSE_SIZE (*osc) > - = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), > - size_one_node); > - if (ptr) > + if (!attach) > + OMP_CLAUSE_SIZE (*osc) > + = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc), > + size_one_node); > + if (ptr || attach_detach) > { > tree cl = insert_struct_comp_map (code, c, NULL, > *prev_list_p, scp); > @@ -9033,11 +9186,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > } > if (!remove > && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER > + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH > + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET > && OMP_CLAUSE_CHAIN (c) > && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP > - && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) > - == GOMP_MAP_ALWAYS_POINTER)) > + && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) > + == GOMP_MAP_ALWAYS_POINTER) > + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) > + == GOMP_MAP_ATTACH_DETACH) > + || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) > + == GOMP_MAP_TO_PSET))) > prev_list_p = list_p; > + > break; > } > flags = GOVD_MAP | GOVD_EXPLICIT; > @@ -9561,6 +9721,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq > *pre_p, > gimplify_omp_ctxp = ctx; > if (struct_map_to_clause) > delete struct_map_to_clause; > + if (struct_deref_set) > + delete struct_deref_set; > } > > /* Return true if DECL is a candidate for shared to firstprivate > @@ -9708,6 +9870,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void > *data) > return 0; > if ((flags & GOVD_SEEN) == 0) > return 0; > + if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0) > + return 0; > if (flags & GOVD_DEBUG_PRIVATE) > { > gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED); > @@ -12762,8 +12926,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq > *pre_p) > && omp_find_clause (OMP_STANDALONE_CLAUSES (expr), > OMP_CLAUSE_FINALIZE)) > { > - /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize" > - semantics apply to all mappings of this OpenACC directive. */ > + /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and > + GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply > + to all mappings of this OpenACC directive. */ > bool finalize_marked = false; > for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN > (c)) > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP) > @@ -12777,10 +12942,19 @@ gimplify_omp_target_update (tree *expr_p, > gimple_seq *pre_p) > OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); > finalize_marked = true; > break; > + case GOMP_MAP_DETACH: > + OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH); > + finalize_marked = true; > + break; > + case GOMP_MAP_STRUCT: > + case GOMP_MAP_FORCE_PRESENT: > + /* Skip over an initial struct or force_present mapping. */ > + break; > default: > - /* Check consistency: libgomp relies on the very first data > - mapping clause being marked, so make sure we did that before > - any other mapping clauses. */ > + /* Check consistency: libgomp relies on the very first > + non-struct, non-force-present data mapping clause being > + marked, so make sure we did that before any other mapping > + clauses. */ > gcc_assert (finalize_marked); > break; > } > diff --git a/gcc/omp-low.c b/gcc/omp-low.c > index d422c205836..3eb7815449a 100644 > --- a/gcc/omp-low.c > +++ b/gcc/omp-low.c > @@ -11439,6 +11439,9 @@ 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; > default: > diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c > index 1cf7a912133..379858d0f1f 100644 > --- a/gcc/tree-pretty-print.c > +++ b/gcc/tree-pretty-print.c > @@ -849,6 +849,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int > spc, dump_flags_t flags) > case GOMP_MAP_LINK: > pp_string (pp, "link"); > break; > + case GOMP_MAP_ATTACH: > + pp_string (pp, "attach"); > + break; > + case GOMP_MAP_DETACH: > + pp_string (pp, "detach"); > + break; > + case GOMP_MAP_FORCE_DETACH: > + pp_string (pp, "force_detach"); > + break; > + case GOMP_MAP_ATTACH_DETACH: > + pp_string (pp, "attach_detach"); > + break; > default: > gcc_unreachable (); > } > @@ -870,6 +882,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int > spc, dump_flags_t flags) > case GOMP_MAP_TO_PSET: > pp_string (pp, " [pointer set, len: "); > break; > + case GOMP_MAP_ATTACH: > + case GOMP_MAP_DETACH: > + case GOMP_MAP_FORCE_DETACH: > + case GOMP_MAP_ATTACH_DETACH: > + pp_string (pp, " [bias: "); > + break; > default: > pp_string (pp, " [len: "); > break; > diff --git a/include/gomp-constants.h b/include/gomp-constants.h > index e8bd52e81bd..f40d6069582 100644 > --- a/include/gomp-constants.h > +++ b/include/gomp-constants.h > @@ -142,7 +142,11 @@ enum gomp_map_kind > /* Do not map, but pointer assign a pointer instead. */ > GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1), > /* Do not map, but pointer assign a reference instead. */ > - GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2) > + GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2), > + /* An attach or detach operation. Rewritten to the appropriate type > during > + gimplification, depending on directive (i.e. "enter data" or > + parallel/kernels region vs. "exit data"). */ > + GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3) > }; > > #define GOMP_MAP_COPY_TO_P(X) \
From 0c1f5b1c22e0c0c3dd0b93697de2235af7e4adfa Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Sat, 21 Dec 2019 21:32:36 +0000
Subject: [PATCH] [PR93026, PR92929] Adjust 'gfortran.dg/goacc/finalize-1.f'
for r279626 changes
gcc/testsuite/
PR fortran/93026
PR middle-end/92929
* gfortran.dg/goacc/finalize-1.f: Adjust.
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@279700 138bc75d-0d04-0410-961f-82ee72b054a4
---
gcc/testsuite/ChangeLog | 6 ++++++
gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 ++--
2 files changed, 8 insertions(+), 2 deletions(-)
diff --git a/gcc/testsuite/ChangeLog b/gcc/testsuite/ChangeLog
index bbd9131e5cd..219ff3b9284 100644
--- a/gcc/testsuite/ChangeLog
+++ b/gcc/testsuite/ChangeLog
@@ -1,3 +1,9 @@
+2019-12-21 Thomas Schwinge <tho...@codesourcery.com>
+
+ PR fortran/93026
+ PR middle-end/92929
+ * gfortran.dg/goacc/finalize-1.f: Adjust.
+
2019-12-21 Harald Anlauf <anl...@gmx.de>
PR fortran/91661
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index ca642156e9f..1e2e3e94b8a 100644
--- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
+++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
@@ -21,7 +21,7 @@
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
@@ -33,5 +33,5 @@
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
--
2.17.1
signature.asc
Description: PGP signature
