@@ -5970,8 +5970,13 @@ omp_notice_variable (struct gimplify_omp
else if (is_scalar)
nflags |= GOVD_FIRSTPRIVATE;
}
+ tree type = TREE_TYPE (decl);
if (nflags == flags
- && !lang_hooks.types.omp_mappable_type (TREE_TYPE (decl)))
+ && gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ type = TREE_TYPE (type);
+ if (nflags == flags
+ && !lang_hooks.types.omp_mappable_type (type))
{
error ("%qD referenced in target region does not have "
"a mappable type", decl);
@@ -6226,7 +6231,7 @@ gimplify_scan_omp_clauses (tree *list_p,
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree, tree> *struct_map_to_clause = NULL;
- tree *orig_list_p = list_p;
+ tree *prev_list_p = NULL;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6506,7 +6511,9 @@ gimplify_scan_omp_clauses (tree *list_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
/* For target {,enter ,exit }data only the array slice is
mapped, but not the pointer to it. */
remove = true;
@@ -6525,7 +6532,9 @@ gimplify_scan_omp_clauses (tree *list_p,
remove = true;
break;
}
- else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
&& TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST)
{
OMP_CLAUSE_SIZE (c)
@@ -6584,6 +6593,25 @@ gimplify_scan_omp_clauses (tree *list_p,
break;
}
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ /* Error recovery. */
+ if (prev_list_p == NULL)
+ {
+ remove = true;
+ break;
+ }
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree ch = OMP_CLAUSE_CHAIN (*prev_list_p);
+ if (ch == NULL_TREE || OMP_CLAUSE_CHAIN (ch) != c)
+ {
+ remove = true;
+ break;
+ }
+ }
+ }
+
tree offset;
HOST_WIDE_INT bitsize, bitpos;
machine_mode mode;
@@ -6603,56 +6631,64 @@ gimplify_scan_omp_clauses (tree *list_p,
splay_tree_node n
= splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
bool ptr = (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER);
- if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
- : GOVD_MAP)) == 0)
+ == GOMP_MAP_ALWAYS_POINTER);
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
{
+ tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+ OMP_CLAUSE_DECL (l) = decl;
+ OMP_CLAUSE_SIZE (l) = size_int (1);
+ if (struct_map_to_clause == NULL)
+ struct_map_to_clause = new hash_map<tree, tree>;
+ struct_map_to_clause->put (decl, l);
if (ptr)
{
+ enum gomp_map_kind mkind
+ = code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_PRIVATE);
- OMP_CLAUSE_DECL (c2) = decl;
- OMP_CLAUSE_CHAIN (c2) = *orig_list_p;
- *orig_list_p = c2;
- if (struct_map_to_clause == NULL)
- struct_map_to_clause = new hash_map<tree, tree>;
- tree *osc;
- if (n == NULL || (n->value & GOVD_MAP) == 0)
- osc = NULL;
- else
- osc = struct_map_to_clause->get (decl);
- if (osc == NULL)
- struct_map_to_clause->put (decl,
- tree_cons (NULL_TREE,
- c,
- NULL_TREE));
- else
- *osc = tree_cons (*osc, c, NULL_TREE);
- flags = GOVD_PRIVATE | GOVD_EXPLICIT;
- goto do_add_decl;
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2)
+ = unshare_expr (OMP_CLAUSE_DECL (c));
+ OMP_CLAUSE_CHAIN (c2) = *prev_list_p;
+ OMP_CLAUSE_SIZE (c2)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (l) = c2;
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p);
+ tree c3
+ = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, mkind);
+ OMP_CLAUSE_DECL (c3)
+ = unshare_expr (OMP_CLAUSE_DECL (c4));
+ OMP_CLAUSE_SIZE (c3)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (c3) = *prev_list_p;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ }
+ *prev_list_p = l;
+ prev_list_p = NULL;
+ }
+ else
+ {
+ OMP_CLAUSE_CHAIN (l) = c;
+ *list_p = l;
+ list_p = &OMP_CLAUSE_CHAIN (l);
}
- *list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
- OMP_CLAUSE_DECL (*list_p) = decl;
- OMP_CLAUSE_SIZE (*list_p) = size_int (1);
- OMP_CLAUSE_CHAIN (*list_p) = c;
- if (struct_map_to_clause == NULL)
- struct_map_to_clause = new hash_map<tree, tree>;
- struct_map_to_clause->put (decl, *list_p);
- list_p = &OMP_CLAUSE_CHAIN (*list_p);
flags = GOVD_MAP | GOVD_EXPLICIT;
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)))
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
flags |= GOVD_SEEN;
goto do_add_decl;
}
else
{
tree *osc = struct_map_to_clause->get (decl);
- tree *sc = NULL, *pt = NULL;
- if (!ptr && TREE_CODE (*osc) == TREE_LIST)
- osc = &TREE_PURPOSE (*osc);
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)))
+ tree *sc = NULL, *scp = NULL;
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
n->value |= GOVD_SEEN;
offset_int o1, o2;
if (offset)
@@ -6661,18 +6697,16 @@ gimplify_scan_omp_clauses (tree *list_p,
o1 = 0;
if (bitpos)
o1 = o1 + bitpos / BITS_PER_UNIT;
- if (ptr)
- pt = osc;
- else
- sc = &OMP_CLAUSE_CHAIN (*osc);
- for (; ptr ? (*pt && (sc = &TREE_VALUE (*pt)))
- : *sc != c;
- ptr ? (pt = &TREE_CHAIN (*pt))
- : (sc = &OMP_CLAUSE_CHAIN (*sc)))
- if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF
- && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
- != INDIRECT_REF)
- && TREE_CODE (OMP_CLAUSE_DECL (*sc)) != ARRAY_REF)
+ for (sc = &OMP_CLAUSE_CHAIN (*osc);
+ *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
+ if (ptr && sc == prev_list_p)
+ break;
+ else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != COMPONENT_REF
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != INDIRECT_REF)
+ && (TREE_CODE (OMP_CLAUSE_DECL (*sc))
+ != ARRAY_REF))
break;
else
{
@@ -6701,6 +6735,8 @@ gimplify_scan_omp_clauses (tree *list_p,
&volatilep, false);
if (base != decl)
break;
+ if (scp)
+ continue;
gcc_assert (offset == NULL_TREE
|| TREE_CODE (offset) == INTEGER_CST);
tree d1 = OMP_CLAUSE_DECL (*sc);
@@ -6739,19 +6775,68 @@ gimplify_scan_omp_clauses (tree *list_p,
o2 = o2 + bitpos2 / BITS_PER_UNIT;
if (wi::ltu_p (o1, o2)
|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
- break;
+ {
+ if (ptr)
+ scp = sc;
+ else
+ break;
+ }
}
+ if (remove)
+ break;
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
if (ptr)
{
- if (!remove)
- *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
- break;
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ tree cl = NULL_TREE;
+ enum gomp_map_kind mkind
+ = code == OMP_TARGET_EXIT_DATA
+ ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+ OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
+ OMP_CLAUSE_DECL (c2)
+ = unshare_expr (OMP_CLAUSE_DECL (c));
+ OMP_CLAUSE_CHAIN (c2) = scp ? *scp : *prev_list_p;
+ OMP_CLAUSE_SIZE (c2)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ cl = scp ? *prev_list_p : c2;
+ if (OMP_CLAUSE_CHAIN (*prev_list_p) != c)
+ {
+ tree c4 = OMP_CLAUSE_CHAIN (*prev_list_p);
+ tree c3
+ = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, mkind);
+ OMP_CLAUSE_DECL (c3)
+ = unshare_expr (OMP_CLAUSE_DECL (c4));
+ OMP_CLAUSE_SIZE (c3)
+ = TYPE_SIZE_UNIT (ptr_type_node);
+ OMP_CLAUSE_CHAIN (c3) = *prev_list_p;
+ if (!scp)
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ else
+ cl = c3;
+ }
+ if (scp)
+ *scp = c2;
+ if (sc == prev_list_p)
+ {
+ *sc = cl;
+ prev_list_p = NULL;
+ }
+ else
+ {
+ *prev_list_p = OMP_CLAUSE_CHAIN (c);
+ list_p = prev_list_p;
+ prev_list_p = NULL;
+ OMP_CLAUSE_CHAIN (c) = *sc;
+ *sc = cl;
+ continue;
+ }
}
- if (!remove)
- OMP_CLAUSE_SIZE (*osc)
- = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
- size_one_node);
- if (!remove && *sc != c)
+ else if (*sc != c)
{
*list_p = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = *sc;
@@ -6760,6 +6845,13 @@ gimplify_scan_omp_clauses (tree *list_p,
}
}
}
+ if (!remove
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+ && OMP_CLAUSE_CHAIN (c)
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ALWAYS_POINTER))
+ prev_list_p = list_p;
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -7248,6 +7340,25 @@ gimplify_adjust_omp_clauses_1 (splay_tre
OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
OMP_CLAUSE_CHAIN (clause) = nc;
}
+ else if (gimplify_omp_ctxp->target_firstprivatize_array_bases
+ && lang_hooks.decls.omp_privatize_by_reference (decl))
+ {
+ OMP_CLAUSE_DECL (clause) = build_simple_mem_ref (decl);
+ OMP_CLAUSE_SIZE (clause)
+ = unshare_expr (TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (decl))));
+ struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
+ gimplify_omp_ctxp = ctx->outer_context;
+ gimplify_expr (&OMP_CLAUSE_SIZE (clause),
+ pre_p, NULL, is_gimple_val, fb_rvalue);
+ gimplify_omp_ctxp = ctx;
+ tree nc = build_omp_clause (OMP_CLAUSE_LOCATION (clause),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (nc) = decl;
+ OMP_CLAUSE_SIZE (nc) = size_zero_node;
+ OMP_CLAUSE_SET_MAP_KIND (nc, GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+ OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (clause);
+ OMP_CLAUSE_CHAIN (clause) = nc;
+ }
else
OMP_CLAUSE_SIZE (clause) = DECL_SIZE_UNIT (decl);
}
@@ -7375,6 +7486,12 @@ gimplify_adjust_omp_clauses (gimple_seq
break;
case OMP_CLAUSE_MAP:
+ if (code == OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ remove = true;
+ break;
+ }
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
{
@@ -7425,7 +7542,9 @@ gimplify_adjust_omp_clauses (gimple_seq
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
/* For GOMP_MAP_FORCE_DEVICEPTR, we'll never enter here, because
for these, TREE_CODE (DECL_SIZE (decl)) will always be
@@ -7468,9 +7587,9 @@ gimplify_adjust_omp_clauses (gimple_seq
{
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
OMP_CLAUSE_SIZE (c) = DECL_SIZE_UNIT (decl);
- if ((n->value & GOVD_SEEN)
- && (n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE)))
- OMP_CLAUSE_MAP_PRIVATE (c) = 1;
+ gcc_assert ((n->value & GOVD_SEEN) == 0
+ || ((n->value & (GOVD_PRIVATE | GOVD_FIRSTPRIVATE))
+ == 0));
}
break;
@@ -2083,7 +2083,9 @@ scan_sharing_clauses (tree clauses, omp_
directly. */
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_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)
@@ -2099,7 +2101,9 @@ scan_sharing_clauses (tree clauses, omp_
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_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
{
if (TREE_CODE (decl) == COMPONENT_REF
|| (TREE_CODE (decl) == INDIRECT_REF
@@ -2128,11 +2132,7 @@ scan_sharing_clauses (tree clauses, omp_
gcc_assert (TREE_CODE (decl2) == INDIRECT_REF);
decl2 = TREE_OPERAND (decl2, 0);
gcc_assert (DECL_P (decl2));
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- install_var_field (decl2, true, 11, ctx);
- else
- install_var_field (decl2, true, 3, ctx);
+ install_var_field (decl2, true, 3, ctx);
install_var_local (decl2, ctx);
install_var_local (decl, ctx);
}
@@ -2143,9 +2143,6 @@ scan_sharing_clauses (tree clauses, omp_
&& !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
&& TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
install_var_field (decl, true, 7, ctx);
- else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- install_var_field (decl, true, 11, ctx);
else
install_var_field (decl, true, 3, ctx);
if (is_gimple_omp_offloaded (ctx->stmt))
@@ -2309,7 +2306,9 @@ scan_sharing_clauses (tree clauses, omp_
break;
decl = OMP_CLAUSE_DECL (c);
if (DECL_P (decl)
- && (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && ((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)
@@ -14363,7 +14362,9 @@ lower_omp_target (gimple_stmt_iterator *
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:
break;
case GOMP_MAP_FORCE_ALLOC:
case GOMP_MAP_FORCE_TO:
@@ -14402,7 +14403,8 @@ lower_omp_target (gimple_stmt_iterator *
}
if (offloaded
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
{
if (TREE_CODE (TREE_TYPE (var)) == ARRAY_TYPE)
{
@@ -14421,12 +14423,6 @@ lower_omp_target (gimple_stmt_iterator *
continue;
}
- if (offloaded && OMP_CLAUSE_MAP_PRIVATE (c))
- {
- map_cnt++;
- continue;
- }
-
if (!maybe_lookup_field (var, ctx))
continue;
@@ -14579,7 +14575,9 @@ lower_omp_target (gimple_stmt_iterator *
nc = c;
ovar = OMP_CLAUSE_DECL (c);
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_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
break;
if (!DECL_P (ovar))
{
@@ -14611,14 +14609,7 @@ lower_omp_target (gimple_stmt_iterator *
gcc_assert (DECL_P (ovar2));
ovar = ovar2;
}
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- {
- if (!maybe_lookup_field ((splay_tree_key) &DECL_UID (ovar),
- ctx))
- continue;
- }
- else if (!maybe_lookup_field (ovar, ctx))
+ if (!maybe_lookup_field (ovar, ctx))
continue;
}
@@ -14628,12 +14619,7 @@ lower_omp_target (gimple_stmt_iterator *
if (nc)
{
var = lookup_decl_in_outer_ctx (ovar, ctx);
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_PRIVATE (c))
- x = build_sender_ref ((splay_tree_key) &DECL_UID (ovar),
- ctx);
- else
- x = build_sender_ref (ovar, ctx);
+ x = build_sender_ref (ovar, ctx);
if (maybe_lookup_oacc_reduction (var, ctx))
{
gcc_checking_assert (offloaded
@@ -15117,7 +15103,7 @@ lower_omp_target (gimple_stmt_iterator *
}
break;
}
- /* Handle GOMP_MAP_FIRSTPRIVATE_POINTER in second pass,
+ /* Handle GOMP_MAP_FIRSTPRIVATE_{POINTER,REFERENCE} in second pass,
so that firstprivate vars holding OMP_CLAUSE_SIZE if needed
are already handled. */
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -15127,7 +15113,8 @@ lower_omp_target (gimple_stmt_iterator *
default:
break;
case OMP_CLAUSE_MAP:
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
HOST_WIDE_INT offset = 0;
@@ -15181,6 +15168,8 @@ lower_omp_target (gimple_stmt_iterator *
}
else
is_ref = is_reference (var);
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ is_ref = false;
bool ref_to_array = false;
if (is_ref)
{
@@ -15232,8 +15221,10 @@ lower_omp_target (gimple_stmt_iterator *
else if (OMP_CLAUSE_CHAIN (c)
&& OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c))
== OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
prev = c;
break;
case OMP_CLAUSE_PRIVATE:
@@ -660,9 +660,15 @@ dump_omp_clause (pretty_printer *pp, tre
case GOMP_MAP_FIRSTPRIVATE_POINTER:
pp_string (pp, "firstprivate");
break;
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ pp_string (pp, "firstprivate ref");
+ break;
case GOMP_MAP_STRUCT:
pp_string (pp, "struct");
break;
+ case GOMP_MAP_ALWAYS_POINTER:
+ pp_string (pp, "always_pointer");
+ break;
default:
gcc_unreachable ();
}
@@ -672,16 +678,22 @@ dump_omp_clause (pretty_printer *pp, tre
print_clause_size:
if (OMP_CLAUSE_SIZE (clause))
{
- if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (clause)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- pp_string (pp, " [pointer assign, bias: ");
- else if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (clause) == GOMP_MAP_TO_PSET)
- pp_string (pp, " [pointer set, len: ");
- else
- pp_string (pp, " [len: ");
+ switch (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP
+ ? OMP_CLAUSE_MAP_KIND (clause) : GOMP_MAP_TO)
+ {
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
+ pp_string (pp, " [pointer assign, bias: ");
+ break;
+ case GOMP_MAP_TO_PSET:
+ pp_string (pp, " [pointer set, len: ");
+ break;
+ default:
+ pp_string (pp, " [len: ");
+ break;
+ }
dump_generic_node (pp, OMP_CLAUSE_SIZE (clause),
spc, flags, false);
pp_right_bracket (pp);
@@ -2902,6 +2902,9 @@ vectorizable_simd_clone_call (gimple *st
case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_CONSTANT_STEP:
case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_CONSTANT_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
/* FORNOW */
i = -1;
break;
@@ -3174,6 +3177,9 @@ vectorizable_simd_clone_call (gimple *st
}
break;
case SIMD_CLONE_ARG_TYPE_LINEAR_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_REF_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_VAL_VARIABLE_STEP:
+ case SIMD_CLONE_ARG_TYPE_LINEAR_UVAL_VARIABLE_STEP:
default:
gcc_unreachable ();
}
@@ -14860,6 +14860,7 @@ c_parser_omp_target_data (location_t loc
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -14993,6 +14994,7 @@ c_parser_omp_target_enter_data (location
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -15079,6 +15081,7 @@ c_parser_omp_target_exit_data (location_
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -15298,6 +15301,7 @@ check_clauses:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -12168,10 +12168,14 @@ handle_omp_array_sections (tree c, bool
break;
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, is_omp
- ? GOMP_MAP_FIRSTPRIVATE_POINTER
- : GOMP_MAP_POINTER);
- if (!is_omp && !c_mark_addressable (t))
+ if (!is_omp)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ else if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && !c_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -12239,7 +12243,7 @@ tree
c_finish_omp_clauses (tree clauses, bool is_omp, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
+ bitmap_head aligned_head, map_head, map_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12256,7 +12260,6 @@ c_finish_omp_clauses (tree clauses, bool
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -12583,6 +12586,12 @@ c_finish_omp_clauses (tree clauses, bool
"%qE appears more than once in data clauses", t);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ && bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&generic_head, DECL_UID (t));
break;
@@ -12604,6 +12613,11 @@ c_finish_omp_clauses (tree clauses, bool
"%qE appears more than once in data clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&firstprivate_head, DECL_UID (t));
break;
@@ -12795,14 +12809,7 @@ c_finish_omp_clauses (tree clauses, bool
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- {
- if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
- break;
- }
- else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
break;
}
}
@@ -12845,13 +12852,13 @@ c_finish_omp_clauses (tree clauses, bool
error ("%qD appears more than once in data clauses", t);
remove = true;
}
- else
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- bitmap_set_bit (&generic_head, DECL_UID (t));
- if (t != OMP_CLAUSE_DECL (c)
- && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
- bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
}
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
@@ -12861,6 +12868,12 @@ c_finish_omp_clauses (tree clauses, bool
error ("%qD appears more than once in map clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
@@ -33797,6 +33797,8 @@ cp_parser_omp_target_data (cp_parser *pa
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -33888,6 +33890,8 @@ cp_parser_omp_target_enter_data (cp_pars
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -33975,6 +33979,8 @@ cp_parser_omp_target_exit_data (cp_parse
map_seen = 3;
break;
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
map_seen |= 1;
@@ -34238,6 +34244,8 @@ check_clauses:
case GOMP_MAP_ALWAYS_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
break;
default:
error_at (OMP_CLAUSE_LOCATION (*pc),
@@ -4907,9 +4907,20 @@ handle_omp_array_sections (tree c, bool
}
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c2, is_omp ? GOMP_MAP_FIRSTPRIVATE_POINTER
- : GOMP_MAP_POINTER);
- if (!is_omp && !cxx_mark_addressable (t))
+ if (!is_omp)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
+ else if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ {
+ t = TREE_OPERAND (t, 0);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && !cxx_mark_addressable (t))
return false;
OMP_CLAUSE_DECL (c2) = t;
t = build_fold_addr_expr (first);
@@ -4927,15 +4938,18 @@ handle_omp_array_sections (tree c, bool
OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
OMP_CLAUSE_CHAIN (c) = c2;
ptr = OMP_CLAUSE_DECL (c2);
- if (!is_omp
+ if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
&& TREE_CODE (TREE_TYPE (ptr)) == REFERENCE_TYPE
&& POINTER_TYPE_P (TREE_TYPE (TREE_TYPE (ptr))))
{
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (c3, OMP_CLAUSE_MAP_KIND (c2));
OMP_CLAUSE_DECL (c3) = ptr;
- OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
+ if (OMP_CLAUSE_MAP_KIND (c2) == GOMP_MAP_ALWAYS_POINTER)
+ OMP_CLAUSE_DECL (c2) = build_simple_mem_ref (ptr);
+ else
+ OMP_CLAUSE_DECL (c2) = convert_from_reference (ptr);
OMP_CLAUSE_SIZE (c3) = size_zero_node;
OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (c2);
OMP_CLAUSE_CHAIN (c2) = c3;
@@ -5659,7 +5673,7 @@ tree
finish_omp_clauses (tree clauses, bool allow_fields, bool declare_simd)
{
bitmap_head generic_head, firstprivate_head, lastprivate_head;
- bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
+ bitmap_head aligned_head, map_head, map_field_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5673,7 +5687,6 @@ finish_omp_clauses (tree clauses, bool a
bitmap_initialize (&aligned_head, &bitmap_default_obstack);
bitmap_initialize (&map_head, &bitmap_default_obstack);
bitmap_initialize (&map_field_head, &bitmap_default_obstack);
- bitmap_initialize (&generic_field_head, &bitmap_default_obstack);
for (pc = &clauses, c = clauses; c ; c = *pc)
{
@@ -5890,6 +5903,12 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_PRIVATE
+ && bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&generic_head, DECL_UID (t));
if (!field_ok)
@@ -5937,6 +5956,11 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
bitmap_set_bit (&firstprivate_head, DECL_UID (t));
goto handle_field_decl;
@@ -6422,7 +6446,10 @@ finish_omp_clauses (tree clauses, bool a
}
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
- t = TREE_OPERAND (t, 0);
+ {
+ t = TREE_OPERAND (t, 0);
+ OMP_CLAUSE_DECL (c) = t;
+ }
if (TREE_CODE (t) == COMPONENT_REF
&& allow_fields
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
@@ -6459,15 +6486,8 @@ finish_omp_clauses (tree clauses, bool a
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER))
- {
- if (bitmap_bit_p (&generic_field_head, DECL_UID (t)))
- break;
- }
- else if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
- break;
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ goto handle_map_references;
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -6475,7 +6495,8 @@ finish_omp_clauses (tree clauses, bool a
if (processing_template_decl)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER)
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER))
break;
if (DECL_P (t))
error ("%qD is not a variable in %qs clause", t,
@@ -6527,17 +6548,13 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in data clauses", t);
remove = true;
}
- else
+ else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
- bitmap_set_bit (&generic_head, DECL_UID (t));
- if (t != OMP_CLAUSE_DECL (c)
- && (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF
- || (REFERENCE_REF_P (OMP_CLAUSE_DECL (c))
- && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c),
- 0))
- == COMPONENT_REF))))
- bitmap_set_bit (&generic_field_head, DECL_UID (t));
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
}
+ else
+ bitmap_set_bit (&generic_head, DECL_UID (t));
}
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
@@ -6547,6 +6564,12 @@ finish_omp_clauses (tree clauses, bool a
error ("%qD appears more than once in map clauses", t);
remove = true;
}
+ else if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears both in data and map clauses", t);
+ remove = true;
+ }
else
{
bitmap_set_bit (&map_head, DECL_UID (t));
@@ -6554,6 +6577,45 @@ finish_omp_clauses (tree clauses, bool a
&& TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
bitmap_set_bit (&map_field_head, DECL_UID (t));
}
+ handle_map_references:
+ if (!remove
+ && !processing_template_decl
+ && allow_fields
+ && TREE_CODE (TREE_TYPE (OMP_CLAUSE_DECL (c))) == REFERENCE_TYPE)
+ {
+ t = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ {
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_ALWAYS_POINTER))
+ {
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ if (TREE_CODE (t) == COMPONENT_REF)
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (c2,
+ GOMP_MAP_FIRSTPRIVATE_REFERENCE);
+ OMP_CLAUSE_DECL (c2) = t;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (c);
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (t);
+ if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
+ OMP_CLAUSE_SIZE (c)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (t)));
+ c = c2;
+ }
+ }
break;
case OMP_CLAUSE_TO_DECLARE:
@@ -4,15 +4,15 @@ void bar (int *);
void
foo (int *p, int q, struct S t, int i, int j, int k, int l)
{
- #pragma omp target map (q), firstprivate (q)
+ #pragma omp target map (q), firstprivate (q) /* { dg-error "appears both in data and map clauses" } */
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)
+ #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])
+ #pragma omp target map (p) , map (p[0]) /* { dg-error "appears both in data and map clauses" } */
bar (p);
#pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */
bar (&q);
@@ -24,17 +24,17 @@ foo (int *p, int q, struct S t, int i, i
bar (&t.r);
#pragma omp target map (t.r) map (t.r) /* { dg-error "appears more than once in map clauses" } */
bar (&t.r);
- #pragma omp target firstprivate (t), map (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)
+ #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)
+ #pragma omp target map (t.s[0]) map (t) /* { dg-error "appears more than once in map clauses" } */
bar (t.s);
- #pragma omp target map (t) map(t.s[0])
+ #pragma omp target map (t) map(t.s[0]) /* { dg-error "appears more than once in map clauses" } */
bar (t.s);
- #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in data clauses" } */
+ #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 more than once in data clauses" } */
+ #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" } */
bar (t.s);
@@ -46,8 +46,8 @@ foo (int *p, int q, struct S t, int i, i
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 more than once in map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 49 } */
- #pragma omp target map (t) map (t.r) firstprivate (t) map (t.s[0]) /* { dg-error "appears more than once in map clauses" } */
- bar (t.s); /* { dg-error "appears more than once in data clauses" "" { target *-*-* } 51 } */
+ #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 *-*-* } 51 } */
}
@@ -111,6 +111,11 @@ enum gomp_map_kind
(address of the last adjacent entry plus its size). */
GOMP_MAP_STRUCT = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 0),
+ /* On a location of a pointer/reference that is assumed to be already mapped
+ earlier, store the translated address of the preceeding mapping.
+ No refcount is bumped by this, and the store is done unconditionally. */
+ GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_FLAG_SPECIAL | 1),
/* Forced deallocation of zero length array section. */
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
= (GOMP_MAP_FLAG_SPECIAL_2
@@ -123,7 +128,9 @@ enum gomp_map_kind
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
- GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1)
+ GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
+ /* Do not map, but pointer assign a reference instead. */
+ GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2)
};
#define GOMP_MAP_COPY_TO_P(X) \
@@ -162,7 +162,20 @@ gomp_map_lookup (splay_tree mem_map, spl
return splay_tree_lookup (mem_map, key);
}
-/* Handle the case where gomp_map_lookup found oldn for newn.
+static inline splay_tree_key
+gomp_map_0len_lookup (splay_tree mem_map, splay_tree_key key)
+{
+ if (key->host_start != key->host_end)
+ return splay_tree_lookup (mem_map, key);
+
+ key->host_end++;
+ splay_tree_key n = splay_tree_lookup (mem_map, key);
+ key->host_end--;
+ return n;
+}
+
+/* Handle the case where gomp_map_lookup, splay_tree_lookup or
+ gomp_map_0len_lookup found oldn for newn.
Helper function of gomp_map_vars. */
static inline void
@@ -306,6 +319,26 @@ gomp_map_fields_existing (struct target_
(void *) cur_node.host_end);
}
+static inline uintptr_t
+gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
+{
+ if (tgt->list[i].key != NULL)
+ return tgt->list[i].key->tgt->tgt_start
+ + tgt->list[i].key->tgt_offset
+ + tgt->list[i].offset;
+ if (tgt->list[i].offset == ~(uintptr_t) 0)
+ return (uintptr_t) hostaddrs[i];
+ if (tgt->list[i].offset == ~(uintptr_t) 1)
+ return 0;
+ if (tgt->list[i].offset == ~(uintptr_t) 2)
+ return tgt->list[i + 1].key->tgt->tgt_start
+ + tgt->list[i + 1].key->tgt_offset
+ + tgt->list[i + 1].offset
+ + (uintptr_t) hostaddrs[i]
+ - (uintptr_t) hostaddrs[i + 1];
+ return tgt->tgt_start + tgt->list[i].offset;
+}
+
attribute_hidden struct target_mem_desc *
gomp_map_vars (struct gomp_device_descr *devicep, size_t mapnum,
void **hostaddrs, void **devaddrs, size_t *sizes, void *kinds,
@@ -396,6 +429,13 @@ gomp_map_vars (struct gomp_device_descr
i--;
continue;
}
+ else if ((kind & typemask) == GOMP_MAP_ALWAYS_POINTER)
+ {
+ tgt->list[i].key = NULL;
+ tgt->list[i].offset = ~(uintptr_t) 1;
+ has_firstprivate = true;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -416,7 +456,7 @@ gomp_map_vars (struct gomp_device_descr
splay_tree_key n;
if ((kind & typemask) == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
{
- n = gomp_map_lookup (mem_map, &cur_node);
+ n = gomp_map_0len_lookup (mem_map, &cur_node);
if (!n)
{
tgt->list[i].key = NULL;
@@ -554,6 +594,32 @@ gomp_map_vars (struct gomp_device_descr
sizes, kinds);
i--;
continue;
+ case GOMP_MAP_ALWAYS_POINTER:
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (mem_map, &cur_node);
+ if (n == NULL
+ || n->host_start > cur_node.host_start
+ || n->host_end < cur_node.host_end)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("always pointer not mapped");
+ }
+ if ((get_kind (short_mapkind, kinds, i - 1) & typemask)
+ != GOMP_MAP_ALWAYS_POINTER)
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i - 1);
+ if (cur_node.tgt_offset)
+ cur_node.tgt_offset -= sizes[i];
+ devicep->host2dev_func (devicep->target_id,
+ (void *) (n->tgt->tgt_start
+ + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start),
+ (void *) &cur_node.tgt_offset,
+ sizeof (void *));
+ cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start - n->host_start;
+ continue;
default:
break;
}
@@ -697,26 +763,7 @@ gomp_map_vars (struct gomp_device_descr
{
for (i = 0; i < mapnum; i++)
{
- if (tgt->list[i].key == NULL)
- {
- if (tgt->list[i].offset == ~(uintptr_t) 0)
- cur_node.tgt_offset = (uintptr_t) hostaddrs[i];
- else if (tgt->list[i].offset == ~(uintptr_t) 1)
- cur_node.tgt_offset = 0;
- else if (tgt->list[i].offset == ~(uintptr_t) 2)
- cur_node.tgt_offset = tgt->list[i + 1].key->tgt->tgt_start
- + tgt->list[i + 1].key->tgt_offset
- + tgt->list[i + 1].offset
- + (uintptr_t) hostaddrs[i]
- - (uintptr_t) hostaddrs[i + 1];
- else
- cur_node.tgt_offset = tgt->tgt_start
- + tgt->list[i].offset;
- }
- else
- cur_node.tgt_offset = tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset
- + tgt->list[i].offset;
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
/* FIXME: see above FIXME comment. */
devicep->host2dev_func (devicep->target_id,
(void *) (tgt->tgt_start
@@ -1551,7 +1598,7 @@ gomp_exit_data (struct gomp_device_descr
cur_node.host_end = cur_node.host_start + sizes[i];
splay_tree_key k = (kind == GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
|| kind == GOMP_MAP_ZERO_LEN_ARRAY_SECTION)
- ? gomp_map_lookup (&devicep->mem_map, &cur_node)
+ ? gomp_map_0len_lookup (&devicep->mem_map, &cur_node)
: splay_tree_lookup (&devicep->mem_map, &cur_node);
if (!k)
continue;
@@ -1783,7 +1830,7 @@ omp_target_is_present (void *ptr, int de
cur_node.host_start = (uintptr_t) ptr;
cur_node.host_end = cur_node.host_start;
- splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
+ splay_tree_key n = gomp_map_0len_lookup (mem_map, &cur_node);
int ret = n != NULL;
gomp_mutex_unlock (&devicep->lock);
return ret;
@@ -41,7 +41,7 @@ main ()
if (omp_target_is_present (q, d) != 1
|| omp_target_is_present (&q[32], d) != 1
- || omp_target_is_present (&q[128], d) != 1)
+ || omp_target_is_present (&q[127], d) != 1)
abort ();
if (omp_target_memcpy (p, q, 128 * sizeof (int), sizeof (int), 0,
@@ -37,58 +37,6 @@ foo (int n)
}
if (err)
abort ();
- int on = n;
- #pragma omp target firstprivate (n) map(tofrom: n)
- {
- n++;
- }
- if (on != n)
- abort ();
- #pragma omp target map(tofrom: n) private (n)
- {
- n = 25;
- }
- if (on != n)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target map(to:a) firstprivate (a) map(from:err) private(i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- if (a[i] != 8 * i)
- err = 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target firstprivate (a) map(to:a) map(from:err) private(i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- if (a[i] != 9 * i)
- err = 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- a[i] += i;
- #pragma omp target map(tofrom:a) map(from:err) private(a, i)
- {
- err = 0;
- for (i = 0; i < n; i++)
- a[i] = 7;
- #pragma omp parallel for reduction(|:err)
- for (i = 0; i < n; i++)
- if (a[i] != 7)
- err |= 1;
- }
- if (err)
- abort ();
- for (i = 0; i < n; i++)
- if (a[i] != 10 * i)
- abort ();
}
int
@@ -1,21 +1,29 @@
extern void abort (void);
-void
+__attribute__((noinline, noclone)) void
foo (int *p, int *q, int *r, int n, int m)
{
int i, err, *s = r;
+ int sep = 1;
+ #pragma omp target map(to:sep)
+ sep = 0;
#pragma omp target data map(to:p[0:8])
{
/* For zero length array sections, p points to the start of
- already mapped range, q to the end of it, and r does not point
- to an mapped range. */
+ already mapped range, q to the end of it (with nothing mapped
+ after it), and r does not point to an mapped range. */
#pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -25,9 +33,14 @@ foo (int *p, int *q, int *r, int n, int
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -38,9 +51,14 @@ foo (int *p, int *q, int *r, int n, int
{
err = 0;
for (i = 0; i < 8; i++)
- if (p[i] != i + 1 || q[i - 8] != i + 1)
+ if (p[i] != i + 1)
err = 1;
- if (p + 8 != q || (r != (int *) 0 && r != s))
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
err = 1;
}
if (err)
@@ -69,7 +87,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)
@@ -81,7 +106,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)
@@ -94,7 +126,14 @@ foo (int *p, int *q, int *r, int n, int
for (i = 0; i < 8; i++)
if (p[i] != i + 1)
err = 1;
- if (q[0] != 9 || r != q + 1)
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
err = 1;
}
if (err)
@@ -0,0 +1,112 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S { char p[64]; int a; int b[2]; long c[4]; int *d; char q[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (struct S s)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int sep = 1;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ int err;
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3]) map(to: sep) map(from: err)
+ {
+ err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+ err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+ s.a = 35; s.b[0] = 36; s.b[1] = 37;
+ s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+ sep = 0;
+ }
+ if (err) abort ();
+ err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+ err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+ if (err) abort ();
+ s.a = 50; s.b[0] = 49; s.b[1] = 48;
+ s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)))
+ abort ();
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+ {
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d))
+ abort ();
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+ {
+ err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+ err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+ s.a = 17; s.b[0] = 18; s.b[1] = 19;
+ s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+ }
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3])
+ }
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+ err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+ if (err) abort ();
+ s.a = 33; s.b[0] = 34; s.b[1] = 35;
+ s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d))
+ abort ();
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3]) map(from: err)
+ {
+ err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+ err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+ s.a = 49; s.b[0] = 48; s.b[1] = 47;
+ s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+ }
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d))
+ abort ();
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3])
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+ err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+ if (err) abort ();
+}
+
+int
+main ()
+{
+ int d[3] = { 18, 19, 20 };
+ struct S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, {} };
+ foo (s);
+ return 0;
+}
@@ -0,0 +1,24 @@
+extern void abort (void);
+
+#pragma omp declare target
+int v = 6;
+#pragma omp end declare target
+
+int
+main ()
+{
+ #pragma omp target /* predetermined map(tofrom: v) */
+ v++;
+ #pragma omp target update from (v)
+ if (v != 7)
+ abort ();
+ #pragma omp parallel private (v) num_threads (1)
+ {
+ #pragma omp target /* predetermined firstprivate(v) */
+ v++;
+ }
+ #pragma omp target update from (v)
+ if (v != 7)
+ abort ();
+ return 0;
+}
@@ -0,0 +1,110 @@
+extern "C" void abort ();
+int x;
+
+__attribute__((noinline, noclone)) void
+foo (int &a, int (&b)[10], short &c, long (&d)[5], int n)
+{
+ int err;
+ int &t = x;
+ int y[n + 1];
+ int (&z)[n + 1] = y;
+ for (int i = 0; i < n + 1; i++)
+ z[i] = i + 27;
+ #pragma omp target enter data map (to: z, c) map (alloc: b, t)
+ #pragma omp target update to (b, t)
+ #pragma omp target map (tofrom: a, d) map (from: b, c) map (alloc: t, z) map (from: err)
+ {
+ err = a++ != 7;
+ for (int i = 0; i < 10; i++)
+ {
+ err |= b[i] != 10 - i;
+ b[i] = i - 16;
+ if (i >= 6) continue;
+ err |= z[i] != i + 27;
+ z[i] = 2 * i + 9;
+ if (i == 5) continue;
+ err |= d[i] != 12L + i;
+ d[i] = i + 7;
+ }
+ err |= c != 25;
+ c = 142;
+ err |= t != 8;
+ t = 19;
+ }
+ if (err) abort ();
+ #pragma omp target update from (z, c)
+ #pragma omp target exit data map (from: b, t) map (release: z, c)
+ if (a != 8 || c != 142 || t != 19)
+ abort ();
+ a = 29;
+ c = 149;
+ t = 15;
+ for (int i = 0; i < 10; i++)
+ {
+ if (b[i] != i - 16) abort ();
+ b[i] = i ^ 1;
+ if (i >= 6) continue;
+ if (z[i] != 2 * i + 9) abort ();
+ z[i]++;
+ if (i == 5) continue;
+ if (d[i] != i + 7) abort ();
+ d[i] = 7 - i;
+ }
+ #pragma omp target defaultmap(tofrom: scalar)
+ {
+ err = a++ != 29;
+ for (int i = 0; i < 10; i++)
+ {
+ err |= b[i] != i ^ 1;
+ b[i] = i + 5;
+ if (i >= 6) continue;
+ err |= z[i] != 2 * i + 10;
+ z[i] = 9 - 3 * i;
+ if (i == 5) continue;
+ err |= d[i] != 7L - i;
+ d[i] = i;
+ }
+ err |= c != 149;
+ c = -2;
+ err |= t != 15;
+ t = 155;
+ }
+ if (err || a != 30 || c != -2 || t != 155)
+ abort ();
+ for (int i = 0; i < 10; i++)
+ {
+ if (b[i] != i + 5) abort ();
+ if (i >= 6) continue;
+ if (z[i] != 9 - 3 * i) abort ();
+ z[i]++;
+ if (i == 5) continue;
+ if (d[i] != i) abort ();
+ }
+ #pragma omp target data map (alloc: z)
+ {
+ #pragma omp target update to (z)
+ #pragma omp target map(from: err)
+ {
+ err = 0;
+ for (int i = 0; i < 6; i++)
+ if (z[i] != 10 - 3 * i) err = 1;
+ else z[i] = i;
+ }
+ if (err) abort ();
+ #pragma omp target update from (z)
+ }
+ for (int i = 0; i < 6; i++)
+ if (z[i] != i)
+ abort ();
+}
+
+int
+main ()
+{
+ int a = 7;
+ int b[10] = { 10, 9, 8, 7, 6, 5, 4, 3, 2, 1 };
+ short c = 25;
+ long d[5] = { 12, 13, 14, 15, 16 };
+ x = 8;
+ foo (a, b, c, d, 5);
+}
@@ -0,0 +1,168 @@
+#include <omp.h>
+#include <stdlib.h>
+
+struct S { char p[64]; int a; int b[2]; long c[4]; int *d; unsigned char &e; char (&f)[2]; short (&g)[4]; int *&h; char q[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (S s)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int sep = 1;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ int err;
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ {
+ err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+ err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+ err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+ err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+ s.a = 35; s.b[0] = 36; s.b[1] = 37;
+ s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+ s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+ s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+ sep = 0;
+ }
+ if (err) abort ();
+ err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+ err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+ err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+ err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+ if (err) abort ();
+ s.a = 50; s.b[0] = 49; s.b[1] = 48;
+ s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+ s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+ s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ {
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+ err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+ err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+ err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+ s.a = 17; s.b[0] = 18; s.b[1] = 19;
+ s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+ s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+ s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+ }
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ }
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+ err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+ err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+ err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+ if (err) abort ();
+ s.a = 33; s.b[0] = 34; s.b[1] = 35;
+ s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+ s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+ s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+ err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+ err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+ err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+ s.a = 49; s.b[0] = 48; s.b[1] = 47;
+ s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+ s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+ s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+ }
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+ err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+ err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+ err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+ if (err) abort ();
+}
+
+int
+main ()
+{
+ int d[3] = { 18, 19, 20 };
+ unsigned char e = 21;
+ char f[2] = { 22, 23 };
+ short g[4] = { 24, 25, 26, 27 };
+ int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+ int *h = hb + 1;
+ S s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+ foo (s);
+}
@@ -0,0 +1,170 @@
+#include <omp.h>
+#include <stdlib.h>
+
+template <typename C, typename I, typename L, typename UC, typename SH>
+struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UC &e; C (&f)[2]; SH (&g)[4]; I *&h; C q[64]; };
+
+template <typename C, typename I, typename L, typename UC, typename SH>
+__attribute__((noinline, noclone)) void
+foo (S<C, I, L, UC, SH> s)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int sep = 1;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ int err;
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ {
+ err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+ err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+ err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+ err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+ s.a = 35; s.b[0] = 36; s.b[1] = 37;
+ s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+ s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+ s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+ sep = 0;
+ }
+ if (err) abort ();
+ err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+ err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+ err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+ err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+ if (err) abort ();
+ s.a = 50; s.b[0] = 49; s.b[1] = 48;
+ s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+ s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+ s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ {
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+ err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+ err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+ err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+ s.a = 17; s.b[0] = 18; s.b[1] = 19;
+ s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+ s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+ s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+ }
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ }
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+ err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+ err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+ err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+ if (err) abort ();
+ s.a = 33; s.b[0] = 34; s.b[1] = 35;
+ s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+ s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+ s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+ err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+ err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+ err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+ s.a = 49; s.b[0] = 48; s.b[1] = 47;
+ s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+ s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+ s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+ }
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+ err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+ err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+ err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+ if (err) abort ();
+}
+
+int
+main ()
+{
+ int d[3] = { 18, 19, 20 };
+ unsigned char e = 21;
+ char f[2] = { 22, 23 };
+ short g[4] = { 24, 25, 26, 27 };
+ int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+ int *h = hb + 1;
+ S<char, int, long, unsigned char, short> s = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+ foo (s);
+}
@@ -0,0 +1,173 @@
+#include <omp.h>
+#include <stdlib.h>
+
+template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR>
+struct S { C p[64]; I a; I b[2]; L c[4]; I *d; UCR e; CAR f; SH g; IPR h; C q[64]; };
+
+template <typename C, typename I, typename L, typename UCR, typename CAR, typename SH, typename IPR>
+__attribute__((noinline, noclone)) void
+foo (S<C, I, L, UCR, CAR, SH, IPR> s)
+{
+ int d = omp_get_default_device ();
+ int id = omp_get_initial_device ();
+ int sep = 1;
+
+ if (d < 0 || d >= omp_get_num_devices ())
+ d = id;
+
+ int err;
+ #pragma omp target map(tofrom: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(to: sep) map(from: err)
+ {
+ err = s.a != 11 || s.b[0] != 12 || s.b[1] != 13;
+ err |= s.c[1] != 15 || s.c[2] != 16 || s.d[-2] != 18 || s.d[-1] != 19 || s.d[0] != 20;
+ err |= s.e != 21 || s.f[0] != 22 || s.f[1] != 23 || s.g[1] != 25 || s.g[2] != 26;
+ err |= s.h[2] != 31 || s.h[3] != 32 || s.h[4] != 33;
+ s.a = 35; s.b[0] = 36; s.b[1] = 37;
+ s.c[1] = 38; s.c[2] = 39; s.d[-2] = 40; s.d[-1] = 41; s.d[0] = 42;
+ s.e = 43; s.f[0] = 44; s.f[1] = 45; s.g[1] = 46; s.g[2] = 47;
+ s.h[2] = 48; s.h[3] = 49; s.h[4] = 50;
+ sep = 0;
+ }
+ if (err) abort ();
+ err = s.a != 35 || s.b[0] != 36 || s.b[1] != 37;
+ err |= s.c[1] != 38 || s.c[2] != 39 || s.d[-2] != 40 || s.d[-1] != 41 || s.d[0] != 42;
+ err |= s.e != 43 || s.f[0] != 44 || s.f[1] != 45 || s.g[1] != 46 || s.g[2] != 47;
+ err |= s.h[2] != 48 || s.h[3] != 49 || s.h[4] != 50;
+ if (err) abort ();
+ s.a = 50; s.b[0] = 49; s.b[1] = 48;
+ s.c[1] = 47; s.c[2] = 46; s.d[-2] = 45; s.d[-1] = 44; s.d[0] = 43;
+ s.e = 42; s.f[0] = 41; s.f[1] = 40; s.g[1] = 39; s.g[2] = 38;
+ s.h[2] = 37; s.h[3] = 36; s.h[4] = 35;
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ #pragma omp target data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ {
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target update to(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 50 || s.b[0] != 49 || s.b[1] != 48;
+ err |= s.c[1] != 47 || s.c[2] != 46 || s.d[-2] != 45 || s.d[-1] != 44 || s.d[0] != 43;
+ err |= s.e != 42 || s.f[0] != 41 || s.f[1] != 40 || s.g[1] != 39 || s.g[2] != 38;
+ err |= s.h[2] != 37 || s.h[3] != 36 || s.h[4] != 35;
+ s.a = 17; s.b[0] = 18; s.b[1] = 19;
+ s.c[1] = 20; s.c[2] = 21; s.d[-2] = 22; s.d[-1] = 23; s.d[0] = 24;
+ s.e = 25; s.f[0] = 26; s.f[1] = 27; s.g[1] = 28; s.g[2] = 29;
+ s.h[2] = 30; s.h[3] = 31; s.h[4] = 32;
+ }
+ #pragma omp target update from(s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ }
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 17 || s.b[0] != 18 || s.b[1] != 19;
+ err |= s.c[1] != 20 || s.c[2] != 21 || s.d[-2] != 22 || s.d[-1] != 23 || s.d[0] != 24;
+ err |= s.e != 25 || s.f[0] != 26 || s.f[1] != 27 || s.g[1] != 28 || s.g[2] != 29;
+ err |= s.h[2] != 30 || s.h[3] != 31 || s.h[4] != 32;
+ if (err) abort ();
+ s.a = 33; s.b[0] = 34; s.b[1] = 35;
+ s.c[1] = 36; s.c[2] = 37; s.d[-2] = 38; s.d[-1] = 39; s.d[0] = 40;
+ s.e = 41; s.f[0] = 42; s.f[1] = 43; s.g[1] = 44; s.g[2] = 45;
+ s.h[2] = 46; s.h[3] = 47; s.h[4] = 48;
+ #pragma omp target enter data map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target enter data map(always, to: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ #pragma omp target map(alloc: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3]) map(from: err)
+ {
+ err = s.a != 33 || s.b[0] != 34 || s.b[1] != 35;
+ err |= s.c[1] != 36 || s.c[2] != 37 || s.d[-2] != 38 || s.d[-1] != 39 || s.d[0] != 40;
+ err |= s.e != 41 || s.f[0] != 42 || s.f[1] != 43 || s.g[1] != 44 || s.g[2] != 45;
+ err |= s.h[2] != 46 || s.h[3] != 47 || s.h[4] != 48;
+ s.a = 49; s.b[0] = 48; s.b[1] = 47;
+ s.c[1] = 46; s.c[2] = 45; s.d[-2] = 44; s.d[-1] = 43; s.d[0] = 42;
+ s.e = 31; s.f[0] = 40; s.f[1] = 39; s.g[1] = 38; s.g[2] = 37;
+ s.h[2] = 36; s.h[3] = 35; s.h[4] = 34;
+ }
+ #pragma omp target exit data map(always, from: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (!omp_target_is_present (&s.a, d)
+ || !omp_target_is_present (s.b, d)
+ || !omp_target_is_present (&s.c[1], d)
+ || !omp_target_is_present (s.d, d)
+ || !omp_target_is_present (&s.d[-2], d)
+ || !omp_target_is_present (&s.e, d)
+ || !omp_target_is_present (s.f, d)
+ || !omp_target_is_present (&s.g[1], d)
+ || !omp_target_is_present (&s.h, d)
+ || !omp_target_is_present (&s.h[2], d))
+ abort ();
+ #pragma omp target exit data map(release: s.a, s.b, s.c[1:2], s.d[-2:3], s.e, s.f, s.g[1:2], s.h[2:3])
+ if (sep
+ && (omp_target_is_present (&s.a, d)
+ || omp_target_is_present (s.b, d)
+ || omp_target_is_present (&s.c[1], d)
+ || omp_target_is_present (s.d, d)
+ || omp_target_is_present (&s.d[-2], d)
+ || omp_target_is_present (&s.e, d)
+ || omp_target_is_present (s.f, d)
+ || omp_target_is_present (&s.g[1], d)
+ || omp_target_is_present (&s.h, d)
+ || omp_target_is_present (&s.h[2], d)))
+ abort ();
+ if (err) abort ();
+ err = s.a != 49 || s.b[0] != 48 || s.b[1] != 47;
+ err |= s.c[1] != 46 || s.c[2] != 45 || s.d[-2] != 44 || s.d[-1] != 43 || s.d[0] != 42;
+ err |= s.e != 31 || s.f[0] != 40 || s.f[1] != 39 || s.g[1] != 38 || s.g[2] != 37;
+ err |= s.h[2] != 36 || s.h[3] != 35 || s.h[4] != 34;
+ if (err) abort ();
+}
+
+int
+main ()
+{
+ int d[3] = { 18, 19, 20 };
+ unsigned char e = 21;
+ char f[2] = { 22, 23 };
+ short g[4] = { 24, 25, 26, 27 };
+ int hb[7] = { 28, 29, 30, 31, 32, 33, 34 };
+ int *h = hb + 1;
+ typedef char (&CAR)[2];
+ typedef short (&SH)[4];
+ S<char, int, long, unsigned char &, CAR, SH, int *&> s
+ = { {}, 11, { 12, 13 }, { 14, 15, 16, 17 }, d + 2, e, f, g, h, {} };
+ foo (s);
+}
@@ -0,0 +1,167 @@
+extern "C" void abort ();
+
+__attribute__((noinline, noclone)) void
+foo (int *&p, int *&q, int *&r, int n, int m)
+{
+ int i, err, *s = r;
+ int sep = 1;
+ #pragma omp target map(to:sep)
+ sep = 0;
+ #pragma omp target data map(to:p[0:8])
+ {
+ /* For zero length array sections, p points to the start of
+ already mapped range, q to the end of it (with nothing mapped
+ after it), and r does not point to an mapped range. */
+ #pragma omp target map(alloc:p[:0]) map(to:q[:0]) map(from:r[:0]) private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Implicit mapping of pointers behaves the same way. */
+ #pragma omp target private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* And zero-length array sections, though not known at compile
+ time, behave the same. */
+ #pragma omp target map(p[:n]) map(tofrom:q[:n]) map(alloc:r[:n]) private(i) map(from:err) firstprivate (s)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (sep)
+ {
+ if (q != (int *) 0 || r != (int *) 0)
+ err = 1;
+ }
+ else if (p + 8 != q || r != s)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Non-zero length array sections, though not known at compile,
+ behave differently. */
+ #pragma omp target map(p[:m]) map(tofrom:q[:m]) map(to:r[:m]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r[0] != 10)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ #pragma omp target data map(to:q[0:1])
+ {
+ /* For zero length array sections, p points to the start of
+ already mapped range, q points to the start of another one,
+ and r to the end of the second one. */
+ #pragma omp target map(to:p[:0]) map(from:q[:0]) map(tofrom:r[:0]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Implicit mapping of pointers behaves the same way. */
+ #pragma omp target private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* And zero-length array sections, though not known at compile
+ time, behave the same. */
+ #pragma omp target map(p[:n]) map(alloc:q[:n]) map(from:r[:n]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9)
+ err = 1;
+ else if (sep)
+ {
+ if (r != (int *) 0)
+ err = 1;
+ }
+ else if (r != q + 1)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ /* Non-zero length array sections, though not known at compile,
+ behave differently. */
+ #pragma omp target map(p[:m]) map(alloc:q[:m]) map(tofrom:r[:m]) private(i) map(from:err)
+ {
+ err = 0;
+ for (i = 0; i < 8; i++)
+ if (p[i] != i + 1)
+ err = 1;
+ if (q[0] != 9 || r[0] != 10)
+ err = 1;
+ }
+ if (err)
+ abort ();
+ }
+ }
+}
+
+int
+main ()
+{
+ int a[32], i;
+ for (i = 0; i < 32; i++)
+ a[i] = i;
+ int *p = a + 1, *q = a + 9, *r = a + 10;
+ foo (p, q, r, 0, 1);
+ return 0;
+}
@@ -0,0 +1,59 @@
+extern "C" void abort ();
+struct S { char a[64]; int (&r)[2]; char b[64]; };
+
+__attribute__((noinline, noclone)) void
+foo (S s, int (&t)[3], int z)
+{
+ int err, sep = 1;
+ // Test that implicit mapping of reference to array does NOT
+ // behave like zero length array sections. s.r can't be used
+ // implicitly, as that means implicit mapping of the whole s
+ // and trying to dereference the references in there is unspecified.
+ #pragma omp target map(from: err) map(to: sep)
+ {
+ err = t[0] != 1 || t[1] != 2 || t[2] != 3;
+ sep = 0;
+ }
+ if (err) abort ();
+ // But explicit zero length array section mapping does.
+ #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0])
+ {
+ if (sep)
+ err = s.r != (int *) 0 || t != (int *) 0;
+ else
+ err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+ }
+ if (err) abort ();
+ // Similarly zero length array section, but unknown at compile time.
+ #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z])
+ {
+ if (sep)
+ err = s.r != (int *) 0 || t != (int *) 0;
+ else
+ err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+ }
+ if (err) abort ();
+ #pragma omp target enter data map (to: s.r, t)
+ // But when already mapped, it binds to existing mappings.
+ #pragma omp target map(from: err) map(tofrom: s.r[:0], t[:0])
+ {
+ err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+ sep = 0;
+ }
+ if (err) abort ();
+ #pragma omp target map(from: err) map(tofrom: s.r[:z], t[:z])
+ {
+ err = t[0] != 1 || t[1] != 2 || t[2] != 3 || s.r[0] != 6 || s.r[1] != 7;
+ sep = 0;
+ }
+ if (err) abort ();
+}
+
+int
+main ()
+{
+ int t[3] = { 1, 2, 3 };
+ int r[2] = { 6, 7 };
+ S s = { {}, r, {} };
+ foo (s, t, 0);
+}