On 2020/10/29 7:49 PM, Jakub Jelinek wrote:
On Wed, Oct 28, 2020 at 06:32:21PM +0800, Chung-Lin Tang wrote:
@@ -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.
I don't like that, it goes against many principles, gimplification really
shouldn't leave around non-GIMPLE IL.
If you need to compare same expression or same expression bases later,
perhaps detect the equalities during gimplification before actually gimplifying
the
clauses and ensure they are gimplified to the same expression or are using
same base (e.g. by adding SAVE_EXPRs or TARGET_EXPRs before the
gimplification).
I have moved that same gimplify_expr call down to below the processing block,
and things still work as expected. My aforementioned gimple-scan-test
modifications
have all been reverted, and all original tests still pass correctly.
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/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 29f385c9368..c2500656193 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)
@@ -9220,6 +9344,14 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq
*pre_p,
}
}
}
+
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
+ == GS_ERROR)
+ {
+ remove = true;
+ break;
+ }
+
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
@@ -9236,10 +9368,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/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/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
