This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
[gomp4.5] Tweak OpenMP 4.5 mapping rules
- From: Jakub Jelinek <jakub at redhat dot com>
- To: gcc-patches at gcc dot gnu dot org
- Date: Thu, 5 Nov 2015 11:28:15 +0100
- Subject: [gomp4.5] Tweak OpenMP 4.5 mapping rules
- Authentication-results: sourceware.org; auth=none
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
Hi!
This patch updates the mapping of vars to what has been ratified
in OpenMP 4.5, or where left unspecified, hopefully follows the
discussed intent.
In particular:
1) for C++ references we map what they refer to, and on target
construct privatize the references themselves (but not what
they point to, because that is mapped)
2) same var may not be present in both data sharing and mapping
clauses
3) structure element based array sections (or C++ references)
don't have the structure elements privatized, but mapped with
an always pointer store at the start of the region (except
exit data; and update doesn't touch the structure elements)
4) omp_target_is_present on one past the last element really
is about what mapping starts at that point, so essentially
it is checking if the first byte at the specified address
is mapped
5) zero length array sections pointing to one past the last
element really are about what mapping starts at that point
>From the above, 3) is really not specified in the standard
and just based on the discussions we had, hopefully OpenMP 5.0 will
clarify, and 4)/5) are fuzzy in the standard and also based
on the discussions.
2015-11-05 Jakub Jelinek <jakub@redhat.com>
gcc/
* gimplify.c (omp_notice_variable): For references check
whether what it refers to has mappable type, rather than
the reference itself.
(gimplify_scan_omp_clauses): Add support for
GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER,
remove old handling of structure element based array sections.
(gimplify_adjust_omp_clauses_1): For implicit references to
variables with reference type and when not ref to scalar or
ref to pointer, map what they refer to using tofrom and
use GOMP_MAP_FIRSTPRIVATE_REFERENCE for the reference.
(gimplify_adjust_omp_clauses): Remove GOMP_MAP_ALWAYS_POINTER
from target exit data. Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE.
Drop OMP_CLAUSE_MAP_PRIVATE support.
* omp-low.c (scan_sharing_clauses): Handle
GOMP_MAP_FIRSTPRIVATE_REFERENCE, drop OMP_CLAUSE_MAP_PRIVATE
support.
(lower_omp_target): Handle GOMP_MAP_FIRSTPRIVATE_REFERENCE
and GOMP_MAP_ALWAYS_POINTER. Drop OMP_CLAUSE_MAP_PRIVATE
support.
* tree-pretty-print.c (dump_omp_clause): Handle
GOMP_MAP_FIRSTPRIVATE_REFERENCE and GOMP_MAP_ALWAYS_POINTER.
Simplify.
* tree-vect-stmts.c (vectorizable_simd_clone_call): Add
SIMD_CLONE_ARG_TYPE_LINEAR_{REF,VAL,UVAL}_VARIABLE_STEP
cases.
gcc/c/
* c-parser.c (c_parser_omp_target_data,
c_parser_omp_target_enter_data,
c_parser_omp_target_exit_data, c_parser_omp_target): Allow
GOMP_MAP_ALWAYS_POINTER.
* c-typeck.c (handle_omp_array_sections): For structure element
based array sections use GOMP_MAP_ALWAYS_POINTER instead of
GOMP_MAP_FIRSTPRIVATE_POINTER.
(c_finish_omp_clauses): Drop generic_field_head, structure
elements are now always mapped even as array section bases,
diagnose same var in data sharing and mapping clauses.
gcc/cp/
* parser.c (cp_parser_omp_target_data,
cp_parser_omp_target_enter_data,
cp_parser_omp_target_exit_data, cp_parser_omp_target): Allow
GOMP_MAP_ALWAYS_POINTER and GOMP_MAP_FIRSTPRIVATE_REFERENCE.
* semantics.c (handle_omp_array_sections): For structure element
based array sections use GOMP_MAP_ALWAYS_POINTER instead of
GOMP_MAP_FIRSTPRIVATE_POINTER.
(finish_omp_clauses): Drop generic_field_head, structure
elements are now always mapped even as array section bases,
diagnose same var in data sharing and mapping clauses.
For references map what they refer to using GOMP_MAP_ALWAYS_POINTER
for structure elements and GOMP_MAP_FIRSTPRIVATE_REFERENCE
otherwise.
gcc/testsuite/
* c-c++-common/gomp/clauses-2.c (foo): Adjust for diagnostics
of variables in both data sharing and mapping clauses and for
structure element based array sections being mapped rather than
privatized.
include/
* gomp-constants.h (enum gomp_map_kind): Add
GOMP_MAP_ALWAYS_POINTER and GOMP_MAP_FIRSTPRIVATE_REFERENCE.
libgomp/
* target.c (gomp_map_0len_lookup, gomp_map_val): New inline
functions.
(gomp_map_vars): Handle GOMP_MAP_ALWAYS_POINTER. For
GOMP_MAP_ZERO_LEN_ARRAY_SECTION use gomp_map_0len_lookup.
Use gomp_map_val function.
(gomp_exit_data): For GOMP_MAP_*ZERO_LEN* use
gomp_map_0len_lookup instead of gomp_map_lookup.
(omp_target_is_present): Use gomp_map_0len_lookup instead of
gomp_map_lookup.
* testsuite/libgomp.c/target-12.c (main): Adjust for
omp_target_is_present change for one-past-last element.
* testsuite/libgomp.c/target-17.c (foo): Drop tests where
the same var is both mapped and privatized.
* testsuite/libgomp.c/target-19.c (foo): Adjust for different
handling of zero-length array sections.
* testsuite/libgomp.c/target-29.c: New test.
* testsuite/libgomp.c/target-30.c: New test.
* testsuite/libgomp.c++/target-14.C: New test.
* testsuite/libgomp.c++/target-15.C: New test.
* testsuite/libgomp.c++/target-16.C: New test.
* testsuite/libgomp.c++/target-17.C: New test.
* testsuite/libgomp.c++/target-18.C: New test.
* testsuite/libgomp.c++/target-19.C: New test.
--- gcc/gimplify.c.jj 2015-11-03 09:21:08.773059315 +0100
+++ gcc/gimplify.c 2015-11-05 10:42:35.772592563 +0100
@@ -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;
--- gcc/omp-low.c.jj 2015-11-03 09:21:08.802058898 +0100
+++ gcc/omp-low.c 2015-11-05 10:44:00.003384618 +0100
@@ -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:
--- gcc/tree-pretty-print.c.jj 2015-11-03 09:21:08.799058941 +0100
+++ gcc/tree-pretty-print.c 2015-11-03 11:58:13.867502798 +0100
@@ -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);
--- gcc/tree-vect-stmts.c.jj 2015-10-14 10:25:50.000000000 +0200
+++ gcc/tree-vect-stmts.c 2015-11-05 10:48:18.025684349 +0100
@@ -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 ();
}
--- gcc/c/c-parser.c.jj 2015-11-03 09:21:09.000000000 +0100
+++ gcc/c/c-parser.c 2015-11-04 14:51:56.710012024 +0100
@@ -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),
--- gcc/c/c-typeck.c.jj 2015-11-03 09:21:08.000000000 +0100
+++ gcc/c/c-typeck.c 2015-11-04 15:17:53.109890507 +0100
@@ -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));
--- gcc/cp/parser.c.jj 2015-11-03 09:21:09.205053109 +0100
+++ gcc/cp/parser.c 2015-11-03 13:31:32.449694248 +0100
@@ -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),
--- gcc/cp/semantics.c.jj 2015-11-03 09:21:08.787059114 +0100
+++ gcc/cp/semantics.c 2015-11-03 16:28:29.133531779 +0100
@@ -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:
--- gcc/testsuite/c-c++-common/gomp/clauses-2.c.jj 2015-11-03 09:21:08.726059990 +0100
+++ gcc/testsuite/c-c++-common/gomp/clauses-2.c 2015-11-04 16:52:53.405837507 +0100
@@ -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 } */
}
--- include/gomp-constants.h.jj 2015-10-26 15:38:20.000000000 +0100
+++ include/gomp-constants.h 2015-11-03 10:13:00.621573428 +0100
@@ -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) \
--- libgomp/target.c.jj 2015-11-02 10:44:09.000000000 +0100
+++ libgomp/target.c 2015-11-04 18:46:11.049937173 +0100
@@ -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;
--- libgomp/testsuite/libgomp.c/target-12.c.jj 2015-10-14 10:24:10.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-12.c 2015-11-05 08:57:43.910783553 +0100
@@ -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,
--- libgomp/testsuite/libgomp.c/target-17.c.jj 2015-10-14 10:24:10.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-17.c 2015-11-04 17:20:39.441143671 +0100
@@ -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
--- libgomp/testsuite/libgomp.c/target-19.c.jj 2015-10-14 10:24:10.000000000 +0200
+++ libgomp/testsuite/libgomp.c/target-19.c 2015-11-05 10:07:56.214421909 +0100
@@ -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)
--- libgomp/testsuite/libgomp.c/target-29.c.jj 2015-11-04 16:54:24.544542125 +0100
+++ libgomp/testsuite/libgomp.c/target-29.c 2015-11-04 18:08:41.861051720 +0100
@@ -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;
+}
--- libgomp/testsuite/libgomp.c/target-30.c.jj 2015-11-04 18:17:50.878194390 +0100
+++ libgomp/testsuite/libgomp.c/target-30.c 2015-11-04 18:17:45.914265082 +0100
@@ -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;
+}
--- libgomp/testsuite/libgomp.c++/target-14.C.jj 2015-11-03 10:13:00.620573442 +0100
+++ libgomp/testsuite/libgomp.c++/target-14.C 2015-11-03 10:13:00.620573442 +0100
@@ -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);
+}
--- libgomp/testsuite/libgomp.c++/target-15.C.jj 2015-11-04 16:39:37.472162348 +0100
+++ libgomp/testsuite/libgomp.c++/target-15.C 2015-11-04 17:59:21.475097239 +0100
@@ -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);
+}
--- libgomp/testsuite/libgomp.c++/target-16.C.jj 2015-11-05 09:55:59.081706150 +0100
+++ libgomp/testsuite/libgomp.c++/target-16.C 2015-11-05 09:58:21.448664482 +0100
@@ -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);
+}
--- libgomp/testsuite/libgomp.c++/target-17.C.jj 2015-11-05 09:59:26.662729254 +0100
+++ libgomp/testsuite/libgomp.c++/target-17.C 2015-11-05 10:05:17.628696101 +0100
@@ -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);
+}
--- libgomp/testsuite/libgomp.c++/target-18.C.jj 2015-11-05 10:06:30.699648230 +0100
+++ libgomp/testsuite/libgomp.c++/target-18.C 2015-11-05 10:20:17.084797486 +0100
@@ -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;
+}
--- libgomp/testsuite/libgomp.c++/target-19.C.jj 2015-11-05 10:18:48.964061178 +0100
+++ libgomp/testsuite/libgomp.c++/target-19.C 2015-11-05 10:30:11.145274934 +0100
@@ -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);
+}
Jakub