[PATCH, 2/3, OpenMP] Target mapping changes for OpenMP 5.0, middle-end parts and compiler testcases
Chung-Lin Tang
cltang@codesourcery.com
Wed Oct 28 10:32:21 GMT 2020
On 2020/10/13 9:31 PM, Jakub Jelinek wrote:
>> +/* Implement OpenMP 5.x map ordering rules for target directives. There are
>> + several rules, and with some level of ambiguity, hopefully we can at least
>> + collect the complexity here in one place. */
>> +
>> +static void
>> +omp_target_reorder_clauses (tree *list_p)
>> +{
> So, first of all, are you convinced we can sort just the explicit clauses
> and leave out the (later on) implicitly added ones?
> If it is possible, sure, it will be easier (because we later on need to deal
> with the GOMP_MAP_STRUCT sorting too).
Yeah, there will probably be more cases to handle later, and possibly sinking
the call to omp_target_reorder_clauses till after the main handling in
gimplify_scan_omp_clauses. But the current routine handles a straightforward
set of cases, which can be grown upon later.
>> + vec<tree> clauses = vNULL;
> Isn't this a memory leak? Nothing frees the vector. Perhaps better
> auto_vec<tree, 32> clauses;
>
>> + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
>> + clauses.safe_push (*cp);
> The rest of the function deals only with OMP_CLAUSE_MAP clauses, wouldn't it
> be better to just push to the vec those clauses and keep other clauses just
> in *list_p chain?
>
>> + /* Collect refs to alloc/release/delete maps. */
>> + vec<tree> ard = vNULL;
> Again, auto_vec<tree, 32> ard;
Thanks for catching this. I'm now using auto_vec now.
>> + tree *cp = list_p;
>> + for (unsigned int i = 0; i < clauses.length (); i++)
>> + if (clauses[i])
>> + {
>> + *cp = clauses[i];
>> + cp = &OMP_CLAUSE_CHAIN (clauses[i]);
>> + }
>> + for (unsigned int i = 0; i < ard.length (); i++)
>> + {
>> + *cp = ard[i];
>> + cp = &OMP_CLAUSE_CHAIN (ard[i]);
>> + }
>> + *cp = NULL_TREE;
>> +
>> + /* OpenMP 5.0 requires that pointer variables are mapped before
>> + its use as a base-pointer. */
>> + for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
>> + if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
>> + {
>> + tree decl = OMP_CLAUSE_DECL (*cp);
>> + gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
>> + if ((k == GOMP_MAP_ALLOC
>> + || k == GOMP_MAP_TO
>> + || k == GOMP_MAP_FROM
>> + || k == GOMP_MAP_TOFROM)
> What about the*ALWAYS* kinds?
Adjustment done, plus re-written so only one pass of this checking is done.
>> + && (TREE_CODE (decl) == INDIRECT_REF
>> + || TREE_CODE (decl) == MEM_REF))
>> + {
>> + tree base_ptr = TREE_OPERAND (decl, 0);
>> + STRIP_TYPE_NOPS (base_ptr);
>> + for (tree *cp2 = &OMP_CLAUSE_CHAIN (*cp); *cp2;
>> + cp2 = &OMP_CLAUSE_CHAIN (*cp2))
>> + if (OMP_CLAUSE_CODE (*cp2) == OMP_CLAUSE_MAP)
>> + {
>> + tree decl2 = OMP_CLAUSE_DECL (*cp2);
>> + gomp_map_kind k2 = OMP_CLAUSE_MAP_KIND (*cp2);
>> + if ((k2 == GOMP_MAP_ALLOC
>> + || k2 == GOMP_MAP_TO
>> + || k2 == GOMP_MAP_FROM
>> + || k2 == GOMP_MAP_TOFROM)
> Again.
>
> This is O(n^2) too, but due to the is_or_contains_p I'm not sure
> if we can avoid it. Perhaps sort the clauses by uid of the base expressions
> and deal with those separately. Maybe let's ignore it for now.
I re-wrote most of omp_target_reorder_clauses to be more efficient. The O(n^2)
issues should be fixed now.
>> @@ -8958,25 +9083,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
>> /* 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;
>> - }
> So what gimplifies those now?
They're gimplified somewhere during omp-low now.
(some gimplify scan testcases were adjusted to accommodate this change)
I don't remember the exact case I encountered, but there were some issues with gimplified
expressions inside the map clauses making some later checking more difficult. I haven't seen
any negative effect of this modification so far.
>> + if (! (code == OMP_TARGET
>> + || code == OMP_TARGET_DATA
>> + || code == OMP_TARGET_ENTER_DATA
>> + || code == OMP_TARGET_EXIT_DATA))
>> + {
> Isn't this just if ((region_type & ORT_ACC) == 0) ? Or do we want
> it for target update too? Though, we wouldn't talk about more than once in
> map clauses then because target update doesn't have those.
It's actually "(region_type & ORT_ACC) != 0", which I now use in the patch.
I originally intended to be careful and only pick those four OpenMP target constructs
to conditionalize on, but so far using the above test works without regressions.
Updated patch attached.
Thanks,
Chung-Lin
-------------- next part --------------
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 29f385c9368..d5048d140d8 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8364,6 +8364,113 @@ extract_base_bit_offset (tree base, tree *base_ref, poly_int64 *bitposp,
return base;
}
+/* Returns true if EXPR is or contains (as a sub-component) BASE_PTR. */
+
+static bool
+is_or_contains_p (tree expr, tree base_ptr)
+{
+ while (expr != base_ptr)
+ if (TREE_CODE (base_ptr) == COMPONENT_REF)
+ base_ptr = TREE_OPERAND (base_ptr, 0);
+ else
+ break;
+ return expr == base_ptr;
+}
+
+/* Implement OpenMP 5.x map ordering rules for target directives. There are
+ several rules, and with some level of ambiguity, hopefully we can at least
+ collect the complexity here in one place. */
+
+static void
+omp_target_reorder_clauses (tree *list_p)
+{
+ /* Collect refs to alloc/release/delete maps. */
+ auto_vec<tree, 32> ard;
+ tree *cp = list_p;
+ while (*cp != NULL_TREE)
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_DELETE))
+ {
+ /* Unlink cp and push to ard. */
+ tree c = *cp;
+ tree nc = OMP_CLAUSE_CHAIN (c);
+ *cp = nc;
+ ard.safe_push (c);
+
+ /* Any associated pointer type maps should also move along. */
+ while (*cp != NULL_TREE
+ && OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ATTACH_DETACH
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (*cp) == GOMP_MAP_TO_PSET))
+ {
+ c = *cp;
+ nc = OMP_CLAUSE_CHAIN (c);
+ *cp = nc;
+ ard.safe_push (c);
+ }
+ }
+ else
+ cp = &OMP_CLAUSE_CHAIN (*cp);
+
+ /* Link alloc/release/delete maps to the end of list. */
+ for (unsigned int i = 0; i < ard.length (); i++)
+ {
+ *cp = ard[i];
+ cp = &OMP_CLAUSE_CHAIN (ard[i]);
+ }
+ *cp = NULL_TREE;
+
+ /* OpenMP 5.0 requires that pointer variables are mapped before
+ its use as a base-pointer. */
+ auto_vec<tree *, 32> atf;
+ for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+ if (OMP_CLAUSE_CODE (*cp) == OMP_CLAUSE_MAP)
+ {
+ /* Collect alloc, to, from, to/from clause tree pointers. */
+ gomp_map_kind k = OMP_CLAUSE_MAP_KIND (*cp);
+ if (k == GOMP_MAP_ALLOC
+ || k == GOMP_MAP_TO
+ || k == GOMP_MAP_FROM
+ || k == GOMP_MAP_TOFROM
+ || k == GOMP_MAP_ALWAYS_TO
+ || k == GOMP_MAP_ALWAYS_FROM
+ || k == GOMP_MAP_ALWAYS_TOFROM)
+ atf.safe_push (cp);
+ }
+
+ for (unsigned int i = 0; i < atf.length (); i++)
+ if (atf[i])
+ {
+ tree *cp = atf[i];
+ tree decl = OMP_CLAUSE_DECL (*cp);
+ if (TREE_CODE (decl) == INDIRECT_REF || TREE_CODE (decl) == MEM_REF)
+ {
+ tree base_ptr = TREE_OPERAND (decl, 0);
+ STRIP_TYPE_NOPS (base_ptr);
+ for (unsigned int j = i + 1; j < atf.length (); j++)
+ {
+ tree *cp2 = atf[j];
+ tree decl2 = OMP_CLAUSE_DECL (*cp2);
+ if (is_or_contains_p (decl2, base_ptr))
+ {
+ /* Move *cp2 to before *cp. */
+ tree c = *cp2;
+ *cp2 = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = *cp;
+ *cp = c;
+ atf[j] = NULL;
+ }
+ }
+ }
+ }
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
@@ -8405,6 +8512,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
+ if (code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ omp_target_reorder_clauses (list_p);
+
while ((c = *list_p) != NULL)
{
bool remove = false;
@@ -8845,15 +8958,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
= get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL,
false);
- omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
- GOVD_FIRSTPRIVATE | GOVD_SEEN);
+ if ((region_type & ORT_TARGET) != 0)
+ omp_add_variable (ctx, OMP_CLAUSE_SIZE (c),
+ GOVD_FIRSTPRIVATE | GOVD_SEEN);
}
+
if (!DECL_P (decl))
{
tree d = decl, *pd;
@@ -8878,7 +8994,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
bool indir_p = false;
tree orig_decl = decl;
tree decl_ref = NULL_TREE;
- if ((region_type & ORT_ACC) != 0
+ if ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA)) != 0
&& TREE_CODE (*pd) == COMPONENT_REF
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
&& code != OACC_UPDATE)
@@ -8886,9 +9002,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
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))
+ if (((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == POINTER_TYPE))
{
indir_p = true;
decl = TREE_OPERAND (decl, 0);
@@ -8915,8 +9033,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (decl != orig_decl && DECL_P (decl) && indir_p)
{
- gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
- : GOMP_MAP_ATTACH;
+ gomp_map_kind k
+ = ((code == OACC_EXIT_DATA || code == OMP_TARGET_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. */
@@ -8925,6 +9044,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree next_clause = OMP_CLAUSE_CHAIN (c);
if (k == GOMP_MAP_ATTACH
&& code != OACC_ENTER_DATA
+ && code != OMP_TARGET_ENTER_DATA
&& (!next_clause
|| (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
|| (OMP_CLAUSE_MAP_KIND (next_clause)
@@ -8972,17 +9092,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
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)
&& 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)
+ && code != OACC_UPDATE
+ && code != OMP_TARGET_UPDATE)
{
if (error_operand_p (decl))
{
@@ -9044,15 +9159,19 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
bool has_attachments = false;
/* For OpenACC, pointers in structs should trigger an
attach action. */
- if (attach_detach && (region_type & ORT_ACC) != 0)
+ if (attach_detach
+ && ((region_type & (ORT_ACC | ORT_TARGET | ORT_TARGET_DATA))
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA))
+
{
/* 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;
+ = ((code == OACC_EXIT_DATA || code == OMP_TARGET_EXIT_DATA)
+ ? GOMP_MAP_DETACH : GOMP_MAP_ATTACH);
OMP_CLAUSE_SET_MAP_KIND (c, k);
has_attachments = true;
}
@@ -9148,33 +9267,38 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
if (scp)
continue;
- tree d1 = OMP_CLAUSE_DECL (*sc);
- tree d2 = OMP_CLAUSE_DECL (c);
- while (TREE_CODE (d1) == ARRAY_REF)
- d1 = TREE_OPERAND (d1, 0);
- while (TREE_CODE (d2) == ARRAY_REF)
- d2 = TREE_OPERAND (d2, 0);
- if (TREE_CODE (d1) == INDIRECT_REF)
- d1 = TREE_OPERAND (d1, 0);
- if (TREE_CODE (d2) == INDIRECT_REF)
- d2 = TREE_OPERAND (d2, 0);
- while (TREE_CODE (d1) == COMPONENT_REF)
- if (TREE_CODE (d2) == COMPONENT_REF
- && TREE_OPERAND (d1, 1)
- == TREE_OPERAND (d2, 1))
- {
+ if ((region_type & ORT_ACC) != 0)
+ {
+ /* This duplicate checking code is currently only
+ enabled for OpenACC. */
+ tree d1 = OMP_CLAUSE_DECL (*sc);
+ tree d2 = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (d1) == ARRAY_REF)
d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
d2 = TREE_OPERAND (d2, 0);
- }
- else
- break;
- if (d1 == d2)
- {
- error_at (OMP_CLAUSE_LOCATION (c),
- "%qE appears more than once in map "
- "clauses", OMP_CLAUSE_DECL (c));
- remove = true;
- break;
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ while (TREE_CODE (d1) == COMPONENT_REF)
+ if (TREE_CODE (d2) == COMPONENT_REF
+ && TREE_OPERAND (d1, 1)
+ == TREE_OPERAND (d2, 1))
+ {
+ d1 = TREE_OPERAND (d1, 0);
+ d2 = TREE_OPERAND (d2, 0);
+ }
+ else
+ break;
+ if (d1 == d2)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE appears more than once in map "
+ "clauses", OMP_CLAUSE_DECL (c));
+ remove = true;
+ break;
+ }
}
if (maybe_lt (offset1, offsetn)
|| (known_eq (offset1, offsetn)
@@ -9236,10 +9360,60 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
+ else
+ {
+ /* DECL_P (decl) == true */
+ tree *sc;
+ if (struct_map_to_clause
+ && (sc = struct_map_to_clause->get (decl)) != NULL
+ && OMP_CLAUSE_MAP_KIND (*sc) == GOMP_MAP_STRUCT
+ && decl == OMP_CLAUSE_DECL (*sc))
+ {
+ /* We have found a map of the whole structure after a
+ leading GOMP_MAP_STRUCT has been created, so refill the
+ leading clause into a map of the whole structure
+ variable, and remove the current one.
+ TODO: we should be able to remove some maps of the
+ following structure element maps if they are of
+ compatible TO/FROM/ALLOC type. */
+ OMP_CLAUSE_SET_MAP_KIND (*sc, OMP_CLAUSE_MAP_KIND (c));
+ OMP_CLAUSE_SIZE (*sc) = unshare_expr (OMP_CLAUSE_SIZE (c));
+ remove = true;
+ break;
+ }
+ }
flags = GOVD_MAP | GOVD_EXPLICIT;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TO
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_TOFROM)
flags |= GOVD_MAP_ALWAYS_TO;
+
+ if ((code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA)
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ {
+ for (struct gimplify_omp_ctx *octx = outer_ctx; octx;
+ octx = octx->outer_context)
+ {
+ splay_tree_node n
+ = splay_tree_lookup (octx->variables,
+ (splay_tree_key) OMP_CLAUSE_DECL (c));
+ /* If this is contained in an outer OpenMP region as a
+ firstprivate value, remove the attach/detach. */
+ if (n && (n->value & GOVD_FIRSTPRIVATE))
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ goto do_add;
+ }
+ }
+
+ enum gomp_map_kind map_kind = (code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ OMP_CLAUSE_SET_MAP_KIND (c, map_kind);
+ }
+
goto do_add;
case OMP_CLAUSE_DEPEND:
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 6d0aa8daeb3..c45ee359e60 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -214,6 +214,21 @@ is_oacc_kernels (omp_context *ctx)
== GF_OMP_TARGET_KIND_OACC_KERNELS));
}
+/* Return true if STMT corresponds to an OpenMP target region. */
+static bool
+is_omp_target (gimple *stmt)
+{
+ if (gimple_code (stmt) == GIMPLE_OMP_TARGET)
+ {
+ int kind = gimple_omp_target_kind (stmt);
+ return (kind == GF_OMP_TARGET_KIND_REGION
+ || kind == GF_OMP_TARGET_KIND_DATA
+ || kind == GF_OMP_TARGET_KIND_ENTER_DATA
+ || kind == GF_OMP_TARGET_KIND_EXIT_DATA);
+ }
+ return false;
+}
+
/* If DECL is the artificial dummy VAR_DECL created for non-static
data member privatization, return the underlying "this" parameter,
otherwise return NULL. */
@@ -1346,7 +1361,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& DECL_P (decl)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
- != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH)
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TO
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_FROM
@@ -1367,6 +1384,40 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
break;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && DECL_P (decl)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt))
+ {
+ /* If this is an offloaded region, an attach operation should
+ only exist when the pointer variable is mapped in a prior
+ clause. */
+ if (is_gimple_omp_offloaded (ctx->stmt))
+ gcc_assert
+ (maybe_lookup_decl (decl, ctx)
+ || (is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
+ && lookup_attribute ("omp declare target",
+ DECL_ATTRIBUTES (decl))));
+
+ /* By itself, attach/detach is generated as part of pointer
+ variable mapping and should not create new variables in the
+ offloaded region, however sender refs for it must be created
+ for its address to be passed to the runtime. */
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c),
+ FIELD_DECL, NULL_TREE, ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ /* To not clash with a map of the pointer variable itself,
+ attach/detach maps have their field looked up by the *clause*
+ tree expression, not the decl. */
+ gcc_assert (!splay_tree_lookup (ctx->field_map,
+ (splay_tree_key) c));
+ splay_tree_insert (ctx->field_map, (splay_tree_key) c,
+ (splay_tree_value) field);
+ break;
+ }
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@@ -1606,6 +1657,11 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable)
break;
+ if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (ctx->stmt)
+ && !is_gimple_omp_offloaded (ctx->stmt))
+ break;
if (DECL_P (decl))
{
if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
@@ -11458,6 +11514,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
case GOMP_MAP_STRUCT:
case GOMP_MAP_ALWAYS_POINTER:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
@@ -11468,8 +11526,6 @@ 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;
@@ -11524,6 +11580,16 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
continue;
}
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ gcc_assert (maybe_lookup_field (c, ctx));
+ map_cnt++;
+ continue;
+ }
+
if (!maybe_lookup_field (var, ctx))
continue;
@@ -11756,14 +11822,28 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (!maybe_lookup_field (ovar, ctx))
+ if (!maybe_lookup_field (ovar, ctx)
+ && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)))
continue;
}
talign = TYPE_ALIGN_UNIT (TREE_TYPE (ovar));
if (DECL_P (ovar) && DECL_ALIGN_UNIT (ovar) > talign)
talign = DECL_ALIGN_UNIT (ovar);
- if (nc)
+
+ if (nc
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ && is_omp_target (stmt))
+ {
+ var = lookup_decl_in_outer_ctx (ovar, ctx);
+ x = build_sender_ref (c, ctx);
+ gimplify_assign (x, build_fold_addr_expr (var), &ilist);
+ }
+ else if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
x = build_sender_ref (ovar, ctx);
diff --git a/gcc/testsuite/c-c++-common/goacc/finalize-1.c b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
index 3d64b2e7cb3..679b0505e19 100644
--- a/gcc/testsuite/c-c++-common/goacc/finalize-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/finalize-1.c
@@ -21,7 +21,7 @@ void f ()
#pragma acc exit data finalize delete (del_f_p[2:5])
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) map\\(firstprivate:del_f_p \\\[pointer assign, bias: 2\\\]\\) finalize;$" 1 "original" } }
- { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\[^ \]+ \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:\\*\\(del_f_p \\+ 2\\) \\\[len: 5\\\]\\) finalize$" 1 "gimple" } } */
#pragma 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 @@ void f ()
#pragma acc exit data copyout (cpo_f_p[4:10]) finalize
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize map\\(from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\) map\\(firstprivate:cpo_f_p \\\[pointer assign, bias: 4\\\]\\);$" 1 "original" } }
- { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\[^ \]+ \\\[len: 10\\\]\\)$" 1 "gimple" } } */
+ { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data finalize map\\(force_from:\\*\\(cpo_f_p \\+ 4\\) \\\[len: 10\\\]\\)$" 1 "gimple" } } */
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index 337c1f7cc77..839269eb62b 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -44,7 +44,7 @@ t1 ()
}
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:\*.*z.? .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:\*.*s\.a.? .len: 40.. map.attach:s.a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.attach:s.e .bias: 0.. map.tofrom:s .len: 32" 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .bias: 0.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
index df405e448b2..9f702ba76f2 100644
--- a/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/struct-enter-exit-data-1.c
@@ -20,8 +20,8 @@ test (int *b, int *c, int *e)
struct str s = { .a = 0, .b = b, .c = c, .d = 0, .e = e, .f = 0 };
#pragma acc enter data copyin(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
- /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*[_0-9]+ \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
+ /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(struct:s \[len: 4\]\) map\(to:s.a \[len: [0-9]+\]\) map\(alloc:s.b \[len: [0-9]+\]\) map\(alloc:s.c \[len: [0-9]+\]\) map\(to:s.f \[len: [0-9]+\]\) map\(to:\*.*s\.b.? \[len: [0-9]+\]\) map\(attach:s.b \[bias: 0\]\) map\(to:\*.*s\.c.? \[len: [0-9]+\]\) map\(attach:s.c \[bias: 0\]\)$} gimple } } */
#pragma acc exit data copyout(s.a, s.b[0:N], s.c[0:N] /* , s.d */ /* , s.e[0:N] */, s.f)
- /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*[_0-9]+ \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
+ /* { dg-final { scan-tree-dump {(?n)#pragma omp target oacc_enter_exit_data map\(from:s.a \[len: [0-9]+\]\) map\(release:s.b \[len: [0-9]+\]\) map\(release:s.c \[len: [0-9]+\]\) map\(from:s.f \[len: [0-9]+\]\) map\(from:\*.*s\.b.? \[len: [0-9]+\]\) map\(detach:s.b \[bias: 0\]\) map\(from:\*.*s\.c.? \[len: [0-9]+\]\) map\(detach:s.c \[bias: 0\]\)$} gimple } } */
}
diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
index ded1d74ccde..bbc8fb4e32b 100644
--- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c
@@ -13,35 +13,35 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (p);
#pragma omp target map (p[0]) map (p) /* { dg-error "appears both in data and map clauses" } */
bar (p);
- #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
+ #pragma omp target map (p) , map (p[0])
bar (p);
#pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
#pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */
bar (p);
- #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t) map (t.r)
bar (&t.r);
- #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.r) map (t)
bar (&t.r);
- #pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.r) map (t.r)
bar (&t.r);
#pragma omp target firstprivate (t), map (t.r) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
#pragma omp target map (t.r) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (&t.r);
- #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.s[0]) map (t)
bar (t.s);
- #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t) map(t.s[0])
bar (t.s);
#pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
- #pragma omp target map (t.s[0]) map (t.s[2]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.s[0]) map (t.s[2])
bar (t.s);
- #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.t[0:2]) map (t.t[4:6])
bar (t.t);
- #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more than once in map clauses" } */
+ #pragma omp target map (t.t[i:j]) map (t.t[k:l])
bar (t.t);
#pragma omp target map (t.s[0]) map (t.r)
bar (t.s);
@@ -50,5 +50,5 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
#pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* { dg-error "appears both in data and map clauses" } */
bar (t.s);
#pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears both in data and map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in map clauses" "" { target *-*-* } .-1 } */
+ bar (t.s);
}
diff --git a/gcc/testsuite/c-c++-common/gomp/map-5.c b/gcc/testsuite/c-c++-common/gomp/map-5.c
new file mode 100644
index 00000000000..1d9d9252864
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/map-5.c
@@ -0,0 +1,24 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+void foo (void)
+{
+ /* Basic test to ensure to,from,tofrom is ordered before alloc,release,delete clauses. */
+ int a, b, c;
+ #pragma omp target enter data map(alloc:a) map(to:b) map(alloc:c)
+ #pragma omp target exit data map(from:a) map(release:b) map(from:c)
+
+ #pragma omp target map(alloc:a) map(tofrom:b) map(alloc:c)
+ a = b = c = 1;
+
+ #pragma omp target enter data map(to:a) map(alloc:b) map(to:c)
+ #pragma omp target exit data map(from:a) map(delete:b) map(from:c)
+}
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(release:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target num_teams.* map\\(tofrom:.* map\\(alloc:.* map\\(alloc:.*" "gimple" } } */
+
+/* { dg-final { scan-tree-dump "pragma omp target enter data map\\(to:.* map\\(to:.* map\\(alloc:.*" "gimple" } } */
+/* { dg-final { scan-tree-dump "pragma omp target exit data map\\(from:.* map\\(from:.* map\\(delete:.*" "gimple" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
index 373bdcb2114..c5ac06943eb 100644
--- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
@@ -12,11 +12,11 @@ program att
!$acc enter data attach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
!$acc exit data detach(myvar%arr2, myptr)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } }
! Test valid usage and processing of the finalize clause.
!$acc exit data detach(myvar%arr2, myptr) finalize
@@ -24,6 +24,6 @@ program att
! For array-descriptor detaches, we no longer generate a "release" mapping
! for the pointed-to data for gimplify.c to turn into "delete". Make sure
! the mapping still isn't there.
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } }
end program att
diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
index a7788580819..0ff2e471180 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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) 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: \\(.*int.*\\) parm\\.1\\.data - \\(.*int.*\\) 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:\\*\\(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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
index 73c4f5a87d0..79bab726dea 100644
--- a/gcc/testsuite/gfortran.dg/gomp/map-2.f90
+++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
@@ -2,5 +2,5 @@ type t
integer :: i
end type t
type(t) v
-!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" }
+!$omp target enter data map(to:v%i, v%i)
end
More information about the Gcc-patches
mailing list