These are the middle-end gimplifier and omp-low changes.
Compiler testcases are also included in this patch.
For attach/detach clauses, I'm currently using the clause tree expression
itself as the key for lookup, to solve the "same-decl" problem when
multiple clauses have the same OMP_CLAUSE_DECL. This is just a special
case right now, have yet to see if this can expand to more general use
between all map clauses.
Thanks,
Chung-Lin
gcc/
* gimplify.c (is_or_contains_p): New static helper function.
(omp_target_reorder_clauses): New function.
(gimplify_scan_omp_clauses): Add use of omp_target_reorder_clauses to
reorder clause list according to OpenMP 5.0 rules. Add handling of
GOMP_MAP_ATTACH_DETACH for OpenMP cases.
* omp-low.c (is_omp_target): New static helper function.
(scan_sharing_clauses): Add scan phase handling of
GOMP_MAP_ATTACH/DETACH
for OpenMP cases.
(lower_omp_target): Add lowering handling of GOMP_MAP_ATTACH/DETACH for
OpenMP cases.
gcc/testsuite/
* c-c++-common/goacc/finalize-1.c: Adjust gimple scanning.
* c-c++-common/goacc/mdc-1.c: Likewise.
* c-c++-common/goacc/struct-enter-exit-data-1.c: Likewise.
* gfortran.dg/goacc/attach-descriptor.f90: Likewise.
* gfortran.dg/goacc/finalize-1.f: Likewise.
* c-c++-common/gomp/clauses-2.c: Remove dg-error cases now valid.
* gfortran.dg/gomp/map-2.f90: Likewise.
* c-c++-common/gomp/map-5.c: New testcase.
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 23d0e2511f7..0ad141c5b3f 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -8350,14 +8350,126 @@ extract_base_bit_offset (tree base, tree *base_ref,
poly_int64 *bitposp,
/* Set *BASE_REF if BASE was a dereferenced reference variable. */
if (base_ref && orig_base != base)
*base_ref = orig_base;
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)
+{
+ vec<tree> clauses = vNULL;
+ for (tree *cp = list_p; *cp; cp = &OMP_CLAUSE_CHAIN (*cp))
+ clauses.safe_push (*cp);
+
+ /* Collect refs to alloc/release/delete maps. */
+ vec<tree> ard = vNULL;
+ for (unsigned int i = 0; i < clauses.length (); i++)
+ if (OMP_CLAUSE_CODE (clauses[i]) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_ALLOC
+ || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (clauses[i]) == GOMP_MAP_DELETE))
+ {
+ ard.safe_push (clauses[i]);
+ clauses[i] = NULL_TREE;
+
+ unsigned int j;
+ for (j = i + 1; j < clauses.length (); j++)
+ {
+ /* Any associated pointer type maps should move along. */
+ tree nc = clauses[j];
+ if (OMP_CLAUSE_CODE (nc) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_REFERENCE
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ATTACH_DETACH
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_TO_PSET))
+ {
+ ard.safe_push (nc);
+ clauses[j] = NULL_TREE;
+ }
+ else
+ break;
+ }
+ i = j - 1;
+ }
+
+ 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)
+ && (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)
+ && 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;
+ if (*cp2 == NULL_TREE)
+ break;
+ }
+ }
+ }
+ }
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
static void
gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
enum omp_region_type region_type,
enum tree_code code)
@@ -8391,14 +8503,20 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
case OACC_PARALLEL:
case OACC_KERNELS:
ctx->target_firstprivatize_array_bases = true;
default:
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;
bool notice_outer = true;
const char *check_non_private = NULL;
unsigned int flags;
tree decl;
@@ -8831,23 +8949,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
NULL, is_gimple_val, fb_rvalue) == GS_ERROR)
{
remove = true;
break;
}
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;
if (TREE_CODE (d) == ARRAY_REF)
{
while (TREE_CODE (d) == ARRAY_REF)
d = TREE_OPERAND (d, 0);
@@ -8864,25 +8985,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
{
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
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)
{
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);
}
if (TREE_CODE (decl) == INDIRECT_REF
&& DECL_P (TREE_OPERAND (decl, 0))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
@@ -8901,24 +9024,26 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
&& DECL_P (TREE_OPERAND (decl, 0))
&& (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
== 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;
+ 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. */
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
+ && code != OMP_TARGET_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)))
@@ -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;
- }
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))
{
remove = true;
break;
}
@@ -9030,23 +9150,27 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
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)
+ 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;
}
if (n == NULL || (n->value & GOVD_MAP) == 0)
{
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
@@ -9134,41 +9258,49 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
tree base
= extract_base_bit_offset (sc_decl, NULL,
&bitposn, &offsetn);
if (base != decl)
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 (! (code == OMP_TARGET
+ || code == OMP_TARGET_DATA
+ || code == OMP_TARGET_ENTER_DATA
+ || code == OMP_TARGET_EXIT_DATA))
+ {
+ /* 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)
&& maybe_lt (bitpos1, bitposn)))
{
if (ptr || attach_detach)
scp = sc;
@@ -9222,18 +9354,68 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
== GOMP_MAP_ATTACH_DETACH)
|| (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
== GOMP_MAP_TO_PSET)))
prev_list_p = list_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:
if (OMP_CLAUSE_DEPEND_KIND (c) == OMP_CLAUSE_DEPEND_SINK)
{
tree deps = OMP_CLAUSE_DECL (c);
while (deps && TREE_CODE (deps) == TREE_LIST)
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index 53efe5f750c..8d50774384a 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -210,14 +210,29 @@ is_oacc_kernels (omp_context *ctx)
{
enum gimple_code outer_type = gimple_code (ctx->stmt);
return ((outer_type == GIMPLE_OMP_TARGET)
&& (gimple_omp_target_kind (ctx->stmt)
== 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. */
tree
omp_member_access_dummy_var (tree decl)
{
@@ -1342,15 +1357,17 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
don't need to be copied, the receiver side will use them
directly. However, global variables with "omp declare target link"
attribute need to be copied. Or when ALWAYS modifier is used. */
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& 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
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_TOFROM
&& is_global_var (maybe_lookup_decl_in_outer_ctx (decl, ctx))
&& varpool_node::get_create (decl)->offloadable
&& !lookup_attribute ("omp declare target link",
@@ -1362,14 +1379,48 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
/* Ignore GOMP_MAP_POINTER kind for arrays in regions that are
not offloaded; there is nothing to map for those. */
if (!is_gimple_omp_offloaded (ctx->stmt)
&& !POINTER_TYPE_P (TREE_TYPE (decl))
&& !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)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
{
if (TREE_CODE (decl) == COMPONENT_REF
|| (TREE_CODE (decl) == INDIRECT_REF
@@ -1601,14 +1652,19 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
&& ((OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& (OMP_CLAUSE_MAP_KIND (c)
!= GOMP_MAP_FIRSTPRIVATE_REFERENCE))
|| TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
&& 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
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE
&& !COMPLETE_TYPE_P (TREE_TYPE (decl)))
{
@@ -11405,26 +11461,26 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
case GOMP_MAP_ALWAYS_TO:
case GOMP_MAP_ALWAYS_FROM:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
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:
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FORCE_FROM:
case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_FORCE_PRESENT:
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:
gcc_unreachable ();
}
#endif
@@ -11471,14 +11527,24 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
x = build_simple_mem_ref (x);
SET_DECL_VALUE_EXPR (new_var, x);
DECL_HAS_VALUE_EXPR_P (new_var) = 1;
}
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;
/* Don't remap compute constructs' reduction variables, because the
intermediate result must be local to each gang. */
if (offloaded && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c)))
@@ -11703,22 +11769,36 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
{
tree ovar2 = DECL_VALUE_EXPR (ovar);
gcc_assert (TREE_CODE (ovar2) == INDIRECT_REF);
ovar2 = TREE_OPERAND (ovar2, 0);
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);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
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
@@ -17,21 +17,21 @@ void f ()
#pragma acc exit data finalize delete (del_f)
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_f\\) finalize;$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1
"gimple" } } */
#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" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
*/
#pragma acc exit data copyout (cpo_f) finalize
/* { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data finalize
map\\(from:cpo_f\\);$" 1 "original" } }
{ dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data finalize map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\)$"
1 "gimple" } } */
#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
@@ -40,15 +40,15 @@ t1 ()
#pragma acc exit data detach(a) finalize
#pragma acc exit data detach(s.a) finalize
}
}
/* { 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" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
map.to:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
map.detach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data
map.attach:s.e .bias: 0.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data
map.release:a .len: 8.." 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
@@ -16,12 +16,12 @@ struct str {
void
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
@@ -9,46 +9,46 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l)
bar (&q);
#pragma omp target map (p[0]) firstprivate (p) /* { dg-error "appears more
than once in data clauses" } */
bar (p);
#pragma omp target firstprivate (p), map (p[0]) /* { dg-error "appears more
than once in data clauses" } */
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);
#pragma omp target map (t.r) ,map (t.s[0])
bar (t.s);
#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
@@ -8,22 +8,22 @@ program att
end type t
type(t) :: myvar
integer, target :: tarr(10)
integer, pointer :: myptr(:)
!$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
! { 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\\\]\\) finalize;$" 1 "original" } }
! 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
@@ -17,21 +17,21 @@
!$ACC EXIT DATA FINALIZE DELETE (del_f)
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(release:del_f\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1
"gimple" } }
!$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" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(from:cpo_r \\\[len: \[0-9\]+\\\]\\)$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_f) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data
map\\(from:cpo_f\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target
oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$"
1 "gimple" } }
!$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
@@ -1,6 +1,6 @@
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
