On Fri, Aug 28, 2015 at 08:13:35PM +0200, Jakub Jelinek wrote:
> Here is my current WIP on further structure element mapping support
> (so, structure element {pointer,reference to pointer,reference to array}
> based array sections, start of C++ support (still need to add tests for
> template instantiation and verify it works properly)).
> I have still pending questions on mapping of references (other than
> array sections) and structure element references pending, hope they will be
> responded to soon and will be able to commit this next week.
And here is the version I've committed. The C++ references (other than
array sections) aren't finished, as I haven't heard from omp-lang on this
topic yet.
Also, another known still broken case is zero length array section handling
on target enter data and target exit data constructs (apparently if only
zero length based array section appears in target enter data construct,
then we treat it as if that construct is exit data instead, plus
delete on zero length array sections is broken too).
For delete of zero length array sections we'll need a new map kind
in any case, for enter data vs. exit data distinction perhaps when we add
a flags parameter to hold e.g. the nowait flag, we can add the exit data
flag (vs. enter data) bit there too and stop using the heuristics.
And we are missing a testcase to test private/firstprivate clauses
on target construct with C++ data members (both normal and in template).
2015-08-31 Jakub Jelinek <ja...@redhat.com>
* gimplify.c (gimplify_scan_omp_clauses): Handle
struct element GOMP_MAP_FIRSTPRIVATE_POINTER.
(gimplify_adjust_omp_clauses): Add CODE argument.
Handle removal of GOMP_MAP_FIRSTPRIVATE_POINTER
struct elements for struct not seen in target body.
Handle removal of struct mapping if struct is not
seen in target body. Remove GOMP_MAP_STRUCT
map clause on OMP_TARGET_EXIT_DATA.
(gimplify_oacc_cache, gimplify_omp_parallel, gimplify_omp_task,
gimplify_omp_for, gimplify_omp_workshare, gimplify_omp_target_update,
gimplify_expr): Adjust callers.
* omp-low.c (scan_sharing_clauses): Handle struct element
GOMP_MAP_FIRSTPRIVATE_POINTER.
(lower_omp_target): Likewise.
gcc/c/
* c-typeck.c (handle_omp_array_sections_1): Add IS_OMP argument, pass
it down recursively. Handle struct element based array sections.
(handle_omp_array_sections): Adjust caller. Handle struct element
based array sections.
(c_finish_omp_clauses): Handle struct element based array sections.
Use generic_head instead of map_head for GOMP_MAP_FIRSTPRIVATE_POINTER
duplicate testing.
gcc/cp/
* parser.c (cp_parser_omp_var_list_no_open): Parse struct element
on map/to/from clauses.
(cp_parser_omp_clause_map): Fix up parsing of delete kind.
* pt.c (tsubst_expr): For OMP_TARGET{,_DATA} pass true instead of
false to allows_field.
* semantics.c (handle_omp_array_sections_1): Add IS_OMP argument,
pass it down recursively. Handle struct element based array sections.
(handle_omp_array_sections): Adjust caller. Handle struct element
based array sections.
(finish_omp_clauses): Handle struct element mappings and struct
element based array sections. Use generic_head instead of map_head
for GOMP_MAP_FIRSTPRIVATE_POINTER duplicate testing.
gcc/testsuite/
* c-c++-common/gomp/clauses-2.c: New test.
* c-c++-common/gomp/clauses-3.c: New test.
libgomp/
* target.c (GOMP_target_enter_exit_data): Allow GOMP_MAP_STRUCT
for enter data and handle it properly.
* testsuite/libgomp.c++/target-10.C: New test.
* testsuite/libgomp.c++/target-11.C: New test.
* testsuite/libgomp.c++/target-12.C: New test.
* testsuite/libgomp.c/target-21.c (z): New variable.
(struct S, main): Add tests for struct element array based array
sections.
* testsuite/libgomp.c/target-22.c: New test.
* testsuite/libgomp.c/target-23.c: New test.
--- gcc/gimplify.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/gimplify.c 2015-08-31 14:52:32.804028967 +0200
@@ -6203,6 +6203,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;
ctx = new_omp_context (region_type);
outer_ctx = ctx->outer_context;
@@ -6443,13 +6444,31 @@ gimplify_scan_omp_clauses (tree *list_p,
}
if (!DECL_P (decl))
{
+ tree d = decl, *pd;
+ if (TREE_CODE (d) == ARRAY_REF)
+ {
+ while (TREE_CODE (d) == ARRAY_REF)
+ d = TREE_OPERAND (d, 0);
+ if (TREE_CODE (d) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (d)) == ARRAY_TYPE)
+ decl = d;
+ }
+ pd = &OMP_CLAUSE_DECL (c);
+ if (d == decl
+ && TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ {
+ pd = &TREE_OPERAND (decl, 0);
+ decl = TREE_OPERAND (decl, 0);
+ }
if (TREE_CODE (decl) == COMPONENT_REF)
{
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
}
- if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p,
- NULL, is_gimple_lvalue, fb_lvalue)
+ if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
@@ -6478,18 +6497,49 @@ gimplify_scan_omp_clauses (tree *list_p,
HOST_WIDE_INT bitsize, bitpos;
machine_mode mode;
int unsignedp, volatilep = 0;
- tree base
- = get_inner_reference (OMP_CLAUSE_DECL (c), &bitsize,
- &bitpos, &offset, &mode, &unsignedp,
- &volatilep, false);
+ tree base = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) == INDIRECT_REF)
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &bitsize, &bitpos, &offset,
+ &mode, &unsignedp,
+ &volatilep, false);
gcc_assert (base == decl
&& (offset == NULL_TREE
|| TREE_CODE (offset) == INTEGER_CST));
splay_tree_node n
= splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
- if (n == NULL || (n->value & GOVD_MAP) == 0)
+ bool ptr = (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_POINTER);
+ if (n == NULL || (n->value & (ptr ? GOVD_PRIVATE
+ : GOVD_MAP)) == 0)
{
+ if (ptr)
+ {
+ 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;
+ }
*list_p = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (*list_p, GOMP_MAP_STRUCT);
@@ -6508,6 +6558,9 @@ gimplify_scan_omp_clauses (tree *list_p,
else
{
tree *osc = struct_map_to_clause->get (decl), *sc;
+ tree *pt = NULL;
+ if (!ptr && TREE_CODE (*osc) == TREE_LIST)
+ osc = &TREE_PURPOSE (*osc);
if (OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS)
n->value |= GOVD_SEEN;
offset_int o1, o2;
@@ -6517,25 +6570,58 @@ gimplify_scan_omp_clauses (tree *list_p,
o1 = 0;
if (bitpos)
o1 = o1 + bitpos / BITS_PER_UNIT;
- for (sc = &OMP_CLAUSE_CHAIN (*osc); *sc != c;
- sc = &OMP_CLAUSE_CHAIN (*sc))
- if (TREE_CODE (OMP_CLAUSE_DECL (*sc)) != COMPONENT_REF)
+ 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)
break;
else
{
tree offset2;
HOST_WIDE_INT bitsize2, bitpos2;
- base = get_inner_reference (OMP_CLAUSE_DECL (*sc),
- &bitsize2, &bitpos2,
- &offset2, &mode,
- &unsignedp, &volatilep,
- false);
+ base = OMP_CLAUSE_DECL (*sc);
+ if (TREE_CODE (base) == ARRAY_REF)
+ {
+ while (TREE_CODE (base) == ARRAY_REF)
+ base = TREE_OPERAND (base, 0);
+ if (TREE_CODE (base) != COMPONENT_REF
+ || (TREE_CODE (TREE_TYPE (base))
+ != ARRAY_TYPE))
+ break;
+ }
+ else if (TREE_CODE (base) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (base, 0))
+ == COMPONENT_REF)
+ && (TREE_CODE (TREE_TYPE
+ (TREE_OPERAND (base, 0)))
+ == REFERENCE_TYPE))
+ base = TREE_OPERAND (base, 0);
+ base = get_inner_reference (base, &bitsize2,
+ &bitpos2, &offset2,
+ &mode, &unsignedp,
+ &volatilep, false);
if (base != decl)
break;
gcc_assert (offset == NULL_TREE
|| TREE_CODE (offset) == INTEGER_CST);
tree d1 = OMP_CLAUSE_DECL (*sc);
tree d2 = OMP_CLAUSE_DECL (c);
+ while (TREE_CODE (d1) == ARRAY_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ while (TREE_CODE (d2) == ARRAY_REF)
+ d2 = TREE_OPERAND (d2, 0);
+ if (TREE_CODE (d1) == INDIRECT_REF)
+ d1 = TREE_OPERAND (d1, 0);
+ if (TREE_CODE (d2) == INDIRECT_REF)
+ d2 = TREE_OPERAND (d2, 0);
while (TREE_CODE (d1) == COMPONENT_REF)
if (TREE_CODE (d2) == COMPONENT_REF
&& TREE_OPERAND (d1, 1)
@@ -6564,6 +6650,12 @@ gimplify_scan_omp_clauses (tree *list_p,
|| (wi::eq_p (o1, o2) && bitpos < bitpos2))
break;
}
+ if (ptr)
+ {
+ if (!remove)
+ *pt = tree_cons (TREE_PURPOSE (*osc), c, *pt);
+ break;
+ }
if (!remove)
OMP_CLAUSE_SIZE (*osc)
= size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
@@ -7081,7 +7173,8 @@ gimplify_adjust_omp_clauses_1 (splay_tre
}
static void
-gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p)
+gimplify_adjust_omp_clauses (gimple_seq *pre_p, tree *list_p,
+ enum tree_code code)
{
struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp;
tree c, decl;
@@ -7176,11 +7269,51 @@ gimplify_adjust_omp_clauses (gimple_seq
case OMP_CLAUSE_MAP:
decl = OMP_CLAUSE_DECL (c);
if (!DECL_P (decl))
- break;
+ {
+ if ((ctx->region_type & ORT_TARGET) != 0
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ decl = TREE_OPERAND (decl, 0);
+ if (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 0);
+ if (DECL_P (decl))
+ {
+ n = splay_tree_lookup (ctx->variables,
+ (splay_tree_key) decl);
+ if (!(n->value & GOVD_SEEN))
+ remove = true;
+ }
+ }
+ }
+ break;
+ }
n = splay_tree_lookup (ctx->variables, (splay_tree_key) decl);
if ((ctx->region_type & ORT_TARGET) != 0
&& !(n->value & GOVD_SEEN)
- && !(OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS))
+ && ((OMP_CLAUSE_MAP_KIND (c) & GOMP_MAP_FLAG_ALWAYS) == 0
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT))
+ {
+ remove = true;
+ /* For struct element mapping, if struct is never referenced
+ in target block and none of the mapping has always modifier,
+ remove all the struct element mappings, which immediately
+ follow the GOMP_MAP_STRUCT map clause. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT)
+ {
+ HOST_WIDE_INT cnt = tree_to_shwi (OMP_CLAUSE_SIZE (c));
+ while (cnt--)
+ OMP_CLAUSE_CHAIN (c)
+ = OMP_CLAUSE_CHAIN (OMP_CLAUSE_CHAIN (c));
+ }
+ }
+ else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
+ && code == OMP_TARGET_EXIT_DATA)
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -7337,7 +7470,7 @@ gimplify_oacc_cache (tree *expr_p, gimpl
gimplify_scan_omp_clauses (&OACC_CACHE_CLAUSES (expr), pre_p, ORT_WORKSHARE,
OACC_CACHE);
- gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OACC_CACHE_CLAUSES (expr), OACC_CACHE);
/* TODO: Do something sensible with this information. */
@@ -7369,7 +7502,8 @@ gimplify_omp_parallel (tree *expr_p, gim
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_PARALLEL_CLAUSES (expr),
+ OMP_PARALLEL);
g = gimple_build_omp_parallel (body,
OMP_PARALLEL_CLAUSES (expr),
@@ -7405,7 +7539,7 @@ gimplify_omp_task (tree *expr_p, gimple_
else
pop_gimplify_context (NULL);
- gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_TASK_CLAUSES (expr), OMP_TASK);
g = gimple_build_omp_task (body,
OMP_TASK_CLAUSES (expr),
@@ -7984,7 +8118,8 @@ gimplify_omp_for (tree *expr_p, gimple_s
TREE_OPERAND (TREE_OPERAND (t, 1), 0) = var;
}
- gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_FOR_CLAUSES (orig_for_stmt),
+ TREE_CODE (orig_for_stmt));
int kind;
switch (TREE_CODE (orig_for_stmt))
@@ -8236,7 +8371,7 @@ gimplify_omp_workshare (tree *expr_p, gi
}
else
gimplify_and_add (OMP_BODY (expr), &body);
- gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_CLAUSES (expr), TREE_CODE (expr));
switch (TREE_CODE (expr))
{
@@ -8312,7 +8447,8 @@ gimplify_omp_target_update (tree *expr_p
}
gimplify_scan_omp_clauses (&OMP_STANDALONE_CLAUSES (expr), pre_p,
ORT_WORKSHARE, TREE_CODE (expr));
- gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr));
+ gimplify_adjust_omp_clauses (pre_p, &OMP_STANDALONE_CLAUSES (expr),
+ TREE_CODE (expr));
stmt = gimple_build_omp_target (NULL, kind, OMP_STANDALONE_CLAUSES (expr));
gimplify_seq_add_stmt (pre_p, stmt);
@@ -9396,7 +9532,8 @@ gimplify_expr (tree *expr_p, gimple_seq
gimplify_scan_omp_clauses (&OMP_CRITICAL_CLAUSES (*expr_p),
pre_p, ORT_WORKSHARE, OMP_CRITICAL);
gimplify_adjust_omp_clauses (pre_p,
- &OMP_CRITICAL_CLAUSES (*expr_p));
+ &OMP_CRITICAL_CLAUSES (*expr_p),
+ OMP_CRITICAL);
g = gimple_build_omp_critical (body,
OMP_CRITICAL_NAME (*expr_p),
OMP_CRITICAL_CLAUSES (*expr_p));
--- gcc/omp-low.c.jj 2015-08-24 14:32:06.000000000 +0200
+++ gcc/omp-low.c 2015-08-28 16:51:51.300696145 +0200
@@ -2074,6 +2074,12 @@ scan_sharing_clauses (tree clauses, omp_
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
+ if (TREE_CODE (decl) == COMPONENT_REF
+ || (TREE_CODE (decl) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (decl, 0)) == COMPONENT_REF
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE)))
+ break;
if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST)
{
@@ -13196,7 +13202,9 @@ lower_omp_target (gimple_stmt_iterator *
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
- || !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c))
+ || (!OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c)
+ && (OMP_CLAUSE_MAP_KIND (c)
+ != GOMP_MAP_FIRSTPRIVATE_POINTER)))
map_cnt++;
continue;
}
@@ -13395,6 +13403,9 @@ lower_omp_target (gimple_stmt_iterator *
case OMP_CLAUSE_FROM:
nc = c;
ovar = OMP_CLAUSE_DECL (c);
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ break;
if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13416,10 +13427,6 @@ lower_omp_target (gimple_stmt_iterator *
}
else
{
- if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
- && OMP_CLAUSE_MAP_KIND (c)
- == GOMP_MAP_FIRSTPRIVATE_POINTER)
- break;
if (DECL_SIZE (ovar)
&& TREE_CODE (DECL_SIZE (ovar)) != INTEGER_CST)
{
@@ -13880,10 +13887,19 @@ lower_omp_target (gimple_stmt_iterator *
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
{
location_t clause_loc = OMP_CLAUSE_LOCATION (c);
+ HOST_WIDE_INT offset = 0;
gcc_assert (prev);
var = OMP_CLAUSE_DECL (c);
- if (DECL_SIZE (var)
- && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
+ if (TREE_CODE (var) == INDIRECT_REF
+ && TREE_CODE (TREE_OPERAND (var, 0)) == COMPONENT_REF)
+ var = TREE_OPERAND (var, 0);
+ if (TREE_CODE (var) == COMPONENT_REF)
+ {
+ var = get_addr_base_and_unit_offset (var, &offset);
+ gcc_assert (var != NULL_TREE && DECL_P (var));
+ }
+ else if (DECL_SIZE (var)
+ && TREE_CODE (DECL_SIZE (var)) != INTEGER_CST)
{
tree var2 = DECL_VALUE_EXPR (var);
gcc_assert (TREE_CODE (var2) == INDIRECT_REF);
@@ -13893,7 +13909,29 @@ lower_omp_target (gimple_stmt_iterator *
}
tree new_var = lookup_decl (var, ctx), x;
tree type = TREE_TYPE (new_var);
- bool is_ref = is_reference (var);
+ bool is_ref;
+ if (TREE_CODE (OMP_CLAUSE_DECL (c)) == INDIRECT_REF
+ && (TREE_CODE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0))
+ == COMPONENT_REF))
+ {
+ type = TREE_TYPE (TREE_OPERAND (OMP_CLAUSE_DECL (c), 0));
+ is_ref = true;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else if (TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ {
+ type = TREE_TYPE (OMP_CLAUSE_DECL (c));
+ is_ref = TREE_CODE (type) == REFERENCE_TYPE;
+ new_var = build2 (MEM_REF, type,
+ build_fold_addr_expr (new_var),
+ build_int_cst (build_pointer_type (type),
+ offset));
+ }
+ else
+ is_ref = is_reference (var);
bool ref_to_array = false;
if (is_ref)
{
--- gcc/c/c-typeck.c.jj 2015-07-31 16:58:09.000000000 +0200
+++ gcc/c/c-typeck.c 2015-08-27 18:53:04.122017251 +0200
@@ -11590,13 +11590,39 @@ c_finish_omp_cancellation_point (locatio
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
{
if (error_operand_p (t))
return error_mark_node;
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "bit-field %qE in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
+ while (TREE_CODE (t) == COMPONENT_REF)
+ {
+ if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is a member of a union", t);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (DECL_P (t))
@@ -11617,11 +11643,11 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- return t;
+ return ret;
}
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -11856,7 +11882,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -12027,7 +12054,9 @@ handle_omp_array_sections (tree c, bool
if (size)
size = c_fully_fold (size, false, NULL);
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
gcc_assert (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FORCE_DEVICEPTR);
if (is_omp)
@@ -12118,7 +12147,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;
+ bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
tree c, t, type, *pc;
tree simdlen = NULL_TREE, safelen = NULL_TREE;
bool branch_seen = false;
@@ -12132,6 +12161,7 @@ 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)
{
@@ -12572,6 +12602,31 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
@@ -12614,7 +12669,14 @@ c_finish_omp_clauses (tree clauses, bool
break;
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
- if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ 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;
}
}
@@ -12648,6 +12710,23 @@ c_finish_omp_clauses (tree clauses, bool
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ 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));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
--- gcc/cp/parser.c.jj 2015-07-29 18:52:12.000000000 +0200
+++ gcc/cp/parser.c 2015-08-31 16:04:58.607705130 +0200
@@ -27950,10 +27950,22 @@ cp_parser_omp_var_list_no_open (cp_parse
decl = error_mark_node;
break;
}
- /* FALL THROUGH. */
+ /* FALLTHROUGH. */
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
+ while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+ {
+ location_t loc
+ = cp_lexer_peek_token (parser->lexer)->location;
+ cp_id_kind idk = CP_ID_KIND_NONE;
+ cp_lexer_consume_token (parser->lexer);
+ decl
+ = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+ decl, false,
+ &idk, loc);
+ }
+ /* FALLTHROUGH. */
case OMP_CLAUSE_DEPEND:
case OMP_CLAUSE_REDUCTION:
while (cp_lexer_next_token_is (parser->lexer, CPP_OPEN_SQUARE))
@@ -29655,7 +29667,9 @@ cp_parser_omp_clause_map (cp_parser *par
int nth = 2;
if (cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COMMA)
nth++;
- if (cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
+ if ((cp_lexer_peek_nth_token (parser->lexer, nth)->type == CPP_NAME
+ || (cp_lexer_peek_nth_token (parser->lexer, nth)->keyword
+ == RID_DELETE))
&& (cp_lexer_peek_nth_token (parser->lexer, nth + 1)->type
== CPP_COLON))
{
@@ -29683,8 +29697,6 @@ cp_parser_omp_clause_map (cp_parser *par
kind = always ? GOMP_MAP_ALWAYS_TOFROM : GOMP_MAP_TOFROM;
else if (strcmp ("release", p) == 0)
kind = GOMP_MAP_RELEASE;
- else if (strcmp ("delete", p) == 0)
- kind = GOMP_MAP_DELETE;
else
{
cp_parser_error (parser, "invalid map kind");
@@ -29696,6 +29708,13 @@ cp_parser_omp_clause_map (cp_parser *par
cp_lexer_consume_token (parser->lexer);
cp_lexer_consume_token (parser->lexer);
}
+ else if (cp_lexer_next_token_is_keyword (parser->lexer, RID_DELETE)
+ && cp_lexer_peek_nth_token (parser->lexer, 2)->type == CPP_COLON)
+ {
+ kind = GOMP_MAP_DELETE;
+ cp_lexer_consume_token (parser->lexer);
+ cp_lexer_consume_token (parser->lexer);
+ }
nlist = cp_parser_omp_var_list_no_open (parser, OMP_CLAUSE_MAP, list,
NULL);
--- gcc/cp/pt.c.jj 2015-07-16 17:56:41.000000000 +0200
+++ gcc/cp/pt.c 2015-08-31 11:48:54.628801176 +0200
@@ -14543,7 +14543,7 @@ tsubst_expr (tree t, tree args, tsubst_f
case OMP_TARGET_DATA:
case OMP_TARGET:
- tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, false,
+ tmp = tsubst_omp_clauses (OMP_CLAUSES (t), false, true,
args, complain, in_decl);
keep_next_level (true);
stmt = begin_omp_structured_block ();
@@ -14558,10 +14558,12 @@ tsubst_expr (tree t, tree args, tsubst_f
break;
case OMP_TARGET_UPDATE:
- tmp = tsubst_omp_clauses (OMP_TARGET_UPDATE_CLAUSES (t), false, false,
+ case OMP_TARGET_ENTER_DATA:
+ case OMP_TARGET_EXIT_DATA:
+ tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), false, true,
args, complain, in_decl);
t = copy_node (t);
- OMP_TARGET_UPDATE_CLAUSES (t) = tmp;
+ OMP_STANDALONE_CLAUSES (t) = tmp;
add_stmt (t);
break;
--- gcc/cp/semantics.c.jj 2015-07-31 16:57:22.000000000 +0200
+++ gcc/cp/semantics.c 2015-08-28 19:58:50.108378664 +0200
@@ -4366,7 +4366,8 @@ omp_privatize_field (tree t)
static tree
handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
- bool &maybe_zero_len, unsigned int &first_non_one)
+ bool &maybe_zero_len, unsigned int &first_non_one,
+ bool is_omp)
{
tree ret, low_bound, length, type;
if (TREE_CODE (t) != TREE_LIST)
@@ -4375,6 +4376,34 @@ handle_omp_array_sections_1 (tree c, tre
return error_mark_node;
if (type_dependent_expression_p (t))
return NULL_TREE;
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ ret = t;
+ if (TREE_CODE (t) == COMPONENT_REF
+ && is_omp
+ && (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
+ || OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
+ {
+ if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "bit-field %qE in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ return error_mark_node;
+ }
+ while (TREE_CODE (t) == COMPONENT_REF)
+ {
+ if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0))) == UNION_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is a member of a union", t);
+ return error_mark_node;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ }
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
@@ -4406,15 +4435,15 @@ handle_omp_array_sections_1 (tree c, tre
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
return error_mark_node;
}
- t = convert_from_reference (t);
- return t;
+ ret = convert_from_reference (ret);
+ return ret;
}
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_REDUCTION
&& TREE_CODE (TREE_CHAIN (t)) == FIELD_DECL)
TREE_CHAIN (t) = omp_privatize_field (TREE_CHAIN (t));
ret = handle_omp_array_sections_1 (c, TREE_CHAIN (t), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one, is_omp);
if (ret == error_mark_node || ret == NULL_TREE)
return ret;
@@ -4656,7 +4685,8 @@ handle_omp_array_sections (tree c, bool
unsigned int first_non_one = 0;
auto_vec<tree, 10> types;
tree first = handle_omp_array_sections_1 (c, OMP_CLAUSE_DECL (c), types,
- maybe_zero_len, first_non_one);
+ maybe_zero_len, first_non_one,
+ is_omp);
if (first == error_mark_node)
return true;
if (first == NULL_TREE)
@@ -4824,7 +4854,9 @@ handle_omp_array_sections (tree c, bool
}
OMP_CLAUSE_DECL (c) = first;
OMP_CLAUSE_SIZE (c) = size;
- if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE))
return false;
if (is_omp)
switch (OMP_CLAUSE_MAP_KIND (c))
@@ -5596,7 +5628,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;
+ bitmap_head aligned_head, map_head, map_field_head, generic_field_head;
tree c, t, *pc;
tree safelen = NULL_TREE;
bool branch_seen = false;
@@ -5608,6 +5640,8 @@ finish_omp_clauses (tree clauses, bool a
bitmap_initialize (&lastprivate_head, &bitmap_default_obstack);
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)
{
@@ -6262,12 +6296,90 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ while (TREE_CODE (t) == ARRAY_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_TYPE (t)) == ARRAY_TYPE)
+ {
+ while (TREE_CODE (t) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
+ break;
+ if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ error ("%qD appears more than once in motion"
+ " clauses", t);
+ else
+ error ("%qD appears more than once in map"
+ " clauses", t);
+ remove = true;
+ }
+ else
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
+ }
}
break;
}
if (t == error_mark_node)
- remove = true;
- else if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
+ {
+ remove = true;
+ break;
+ }
+ if (REFERENCE_REF_P (t)
+ && TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
+ t = TREE_OPERAND (t, 0);
+ if (TREE_CODE (t) == COMPONENT_REF
+ && allow_fields
+ && OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
+ {
+ if (type_dependent_expression_p (t))
+ break;
+ if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "bit-field %qE in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ remove = true;
+ }
+ else if (!cp_omp_mappable_type (TREE_TYPE (t)))
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE does not have a mappable type in %qs clause",
+ t, omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ remove = true;
+ }
+ while (TREE_CODE (t) == COMPONENT_REF)
+ {
+ if (TREE_CODE (TREE_TYPE (TREE_OPERAND (t, 0)))
+ == UNION_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%qE is a member of a union", t);
+ remove = true;
+ break;
+ }
+ t = TREE_OPERAND (t, 0);
+ }
+ if (remove)
+ 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 (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
if (processing_template_decl)
break;
@@ -6303,6 +6415,7 @@ finish_omp_clauses (tree clauses, bool a
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_POINTER)))
+ && t == OMP_CLAUSE_DECL (c)
&& !type_dependent_expression_p (t)
&& !cp_omp_mappable_type ((TREE_CODE (TREE_TYPE (t))
== REFERENCE_TYPE)
@@ -6314,6 +6427,27 @@ finish_omp_clauses (tree clauses, bool a
omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
remove = true;
}
+ else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER)
+ {
+ if (bitmap_bit_p (&generic_head, DECL_UID (t))
+ || bitmap_bit_p (&firstprivate_head, DECL_UID (t)))
+ {
+ error ("%qD appears more than once in data clauses", t);
+ remove = true;
+ }
+ else
+ {
+ 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));
+ }
+ }
else if (bitmap_bit_p (&map_head, DECL_UID (t)))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
@@ -6323,7 +6457,12 @@ finish_omp_clauses (tree clauses, bool a
remove = true;
}
else
- bitmap_set_bit (&map_head, DECL_UID (t));
+ {
+ bitmap_set_bit (&map_head, DECL_UID (t));
+ if (t != OMP_CLAUSE_DECL (c)
+ && TREE_CODE (OMP_CLAUSE_DECL (c)) == COMPONENT_REF)
+ bitmap_set_bit (&map_field_head, DECL_UID (t));
+ }
break;
case OMP_CLAUSE_TO_DECLARE:
--- gcc/testsuite/c-c++-common/gomp/clauses-2.c.jj 2015-08-28
10:54:34.545144458 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-2.c 2015-08-28 11:19:58.601066200
+0200
@@ -0,0 +1,53 @@
+struct S { int r; int *s; int t[10]; };
+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)
+ 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)
+ bar (p);
+ #pragma omp target map (p) , map (p[0])
+ bar (p);
+ #pragma omp target map (q) map (q) /* { dg-error "appears more than once in
map clauses" } */
+ bar (&q);
+ #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than
once in data clauses" } */
+ bar (p);
+ #pragma omp target map (t) map (t.r) /* { dg-error "appears more than once
in map clauses" } */
+ bar (&t.r);
+ #pragma omp target map (t.r) map (t) /* { dg-error "appears more than once
in map clauses" } */
+ 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)
+ bar (&t.r);
+ #pragma omp target map (t.r) firstprivate (t)
+ bar (&t.r);
+ #pragma omp target map (t.s[0]) map (t)
+ bar (t.s);
+ #pragma omp target map (t) map(t.s[0])
+ bar (t.s);
+ #pragma omp target firstprivate (t) map (t.s[0]) /* { dg-error "appears more
than once in data clauses" } */
+ bar (t.s);
+ #pragma omp target map (t.s[0]) firstprivate (t) /* { dg-error "appears more
than once in data 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);
+ #pragma omp target map (t.t[0:2]) map (t.t[4:6]) /* { dg-error "appears more
than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.t[i:j]) map (t.t[k:l]) /* { dg-error "appears more
than once in map clauses" } */
+ bar (t.t);
+ #pragma omp target map (t.s[0]) map (t.r)
+ bar (t.s);
+ #pragma omp target map (t.r) ,map (t.s[0])
+ bar (t.s);
+ #pragma omp target map (t.r) map (t) map (t.s[0]) firstprivate (t) /* {
dg-error "appears 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 } */
+}
--- gcc/testsuite/c-c++-common/gomp/clauses-3.c.jj 2015-08-28
19:56:08.924530062 +0200
+++ gcc/testsuite/c-c++-common/gomp/clauses-3.c 2015-08-28 19:48:19.000000000
+0200
@@ -0,0 +1,23 @@
+struct T { int a; int *b; };
+struct S { int *s; char u; struct T v; long x; };
+
+void bar (int *);
+#pragma omp declare target to (bar)
+
+int
+main ()
+{
+ int a[10] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 };
+ struct S s = { a, 5, { 6, a + 5 }, 99L };
+ #pragma omp target map (s.v.a, s.u, s.x)
+ ;
+ #pragma omp target map (s.v.a, s.u, s.x)
+ bar (&s.v.a);
+ #pragma omp target map (s.v.a) map (always, to: s.u) map (s.x)
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ ;
+ #pragma omp target map (s.s[0]) map (s.v.b[:3])
+ bar (s.s);
+ return 0;
+}
--- libgomp/target.c.jj 2015-07-31 16:55:38.000000000 +0200
+++ libgomp/target.c 2015-08-31 15:35:03.670073075 +0200
@@ -1465,7 +1465,8 @@ GOMP_target_enter_exit_data (int device,
if (kind == GOMP_MAP_ALLOC
|| kind == GOMP_MAP_TO
- || kind == GOMP_MAP_ALWAYS_TO)
+ || kind == GOMP_MAP_ALWAYS_TO
+ || kind == GOMP_MAP_STRUCT)
{
is_enter_data = true;
break;
@@ -1483,8 +1484,15 @@ GOMP_target_enter_exit_data (int device,
if (is_enter_data)
for (i = 0; i < mapnum; i++)
- gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
- true, GOMP_MAP_VARS_ENTER_DATA);
+ if ((kinds[i] & 0xff) == GOMP_MAP_STRUCT)
+ {
+ gomp_map_vars (devicep, sizes[i] + 1, &hostaddrs[i], NULL, &sizes[i],
+ &kinds[i], true, GOMP_MAP_VARS_ENTER_DATA);
+ i += sizes[i];
+ }
+ else
+ gomp_map_vars (devicep, 1, &hostaddrs[i], NULL, &sizes[i], &kinds[i],
+ true, GOMP_MAP_VARS_ENTER_DATA);
else
gomp_exit_data (devicep, mapnum, hostaddrs, sizes, kinds);
}
--- libgomp/testsuite/libgomp.c++/target-10.C.jj 2015-08-28
10:57:13.898941691 +0200
+++ libgomp/testsuite/libgomp.c++/target-10.C 2015-08-31 11:06:58.000000000
+0200
@@ -0,0 +1,154 @@
+extern "C" void abort (void);
+union U { int x; long long y; };
+struct T { int a; union U b; int c; };
+struct S { int s; int u; T v; int x[10]; union U w; int y[10]; int z[10]; };
+volatile int z;
+
+template <typename R>
+void
+foo ()
+{
+ R s;
+ s.template s = 0;
+ s.u = 1;
+ s.v.a = 2;
+ s.v.b.y = 3LL;
+ s.v.c = 19;
+ s.w.x = 4;
+ s.template x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
+ int err = 0;
+ #pragma omp target map (to:s.template v.template b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.template z[z + 1:z + 3], err)
+ {
+ err = 0;
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
+ err = 1;
+ s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
+ }
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
+ #pragma omp target data map (tofrom: s)
+ #pragma omp target map (always to: s.template w, s.x[1], err) map
(alloc:s.u, s. template v.template b, s.z[z:z + 1])
+ {
+ err = 0;
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] !=
19)
+ err = 1;
+ s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
+ }
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] !=
23)
+ err = 1;
+ s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
+ }
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
+ abort ();
+}
+
+int
+main ()
+{
+ S s;
+ s.s = 0;
+ s.u = 1;
+ s.v.a = 2;
+ s.v.b.y = 3LL;
+ s.v.c = 19;
+ s.w.x = 4;
+ s.x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
+ int err = 0;
+ #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.z[z + 1:z + 3], err)
+ {
+ err = 0;
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
+ err = 1;
+ s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
+ }
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
+ #pragma omp target data map (tofrom: s)
+ #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b,
s.z[z:z + 1])
+ {
+ err = 0;
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] !=
19)
+ err = 1;
+ s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
+ }
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
+ abort ();
+ s.u++;
+ s.v.a++;
+ s.v.b.y++;
+ s.w.x++;
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
+ {
+ err = 0;
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] !=
23)
+ err = 1;
+ s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
+ }
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
+ abort ();
+ foo <S> ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-11.C.jj 2015-08-28
10:57:16.860900748 +0200
+++ libgomp/testsuite/libgomp.c++/target-11.C 2015-08-31 12:01:17.000000000
+0200
@@ -0,0 +1,121 @@
+extern "C" void abort ();
+struct T { int a; int *b; int c; char (&d)[10]; };
+struct S { int *s; char *u; T v; short *w; short *&x; };
+volatile int z;
+
+template <typename A, typename B, typename C, typename D>
+void
+foo ()
+{
+ A d[10];
+ B *e;
+ C a[32], i;
+ A b[32];
+ B c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ for (i = 0; i < 10; i++)
+ d[i] = 17 + i;
+ e = c + 18;
+ D s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.template u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3], s. template v. template d[z + 1:z +
3]) \
+ map (from: s.w[z:4], s.x[1:3], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 1; i < 4; i++)
+ if (s.v.d[i] != 17 + i)
+ err = 1;
+ else
+ s.v.d[i] = 23 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ for (i = 1; i < 4; i++)
+ s.x[i] = 173 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : ((i >= 19 && i < 22) ? 155 + i
: 64 + i)))
+ abort ();
+ for (i = 0; i < 10; i++)
+ if (d[i] != ((i >= 1 && i < 4) ? 23 + i : 17 + i))
+ abort ();
+}
+
+int
+main ()
+{
+ char d[10];
+ short *e;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ for (i = 0; i < 10; i++)
+ d[i] = 17 + i;
+ e = c + 18;
+ S s = { a, b + 2, { 0, a + 16, 0, d }, c + 3, e };
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3], s.v.d[z + 1:z + 3]) \
+ map (from: s.w[z:4], s.x[1:3], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 1; i < 4; i++)
+ if (s.v.d[i] != 17 + i)
+ err = 1;
+ else
+ s.v.d[i] = 23 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ for (i = 1; i < 4; i++)
+ s.x[i] = 173 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : ((i >= 19 && i < 22) ? 155 + i
: 64 + i)))
+ abort ();
+ for (i = 0; i < 10; i++)
+ if (d[i] != ((i >= 1 && i < 4) ? 23 + i : 17 + i))
+ abort ();
+ foo <char, short, int, S> ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c++/target-12.C.jj 2015-08-31
15:39:10.329714036 +0200
+++ libgomp/testsuite/libgomp.c++/target-12.C 2015-08-31 15:56:32.809545094
+0200
@@ -0,0 +1,93 @@
+extern "C" void abort (void);
+struct S { int s; int *u; int v[5]; };
+volatile int z;
+
+template <typename T>
+void
+foo ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ T s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.template u[0:5]) map (alloc:
s.template v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.template s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] !=
9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+}
+
+int
+main ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] !=
9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (always, delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (always delete : s.v[1:3])
+}
--- libgomp/testsuite/libgomp.c/target-21.c.jj 2015-07-31 17:32:56.000000000
+0200
+++ libgomp/testsuite/libgomp.c/target-21.c 2015-08-31 12:06:13.994068316
+0200
@@ -1,7 +1,8 @@
extern void abort (void);
union U { int x; long long y; };
struct T { int a; union U b; int c; };
-struct S { int s; int u; struct T v; union U w; };
+struct S { int s; int u; struct T v; int x[10]; union U w; int y[10]; int
z[10]; };
+volatile int z;
int
main ()
@@ -13,43 +14,66 @@ main ()
s.v.b.y = 3LL;
s.v.c = 19;
s.w.x = 4;
+ s.x[0] = 7;
+ s.x[1] = 8;
+ s.y[3] = 9;
+ s.y[4] = 10;
+ s.y[5] = 11;
int err = 0;
- #pragma omp target map (to:s.v.b, s.u) map (from: s.w, err)
+ #pragma omp target map (to:s.v.b, s.u, s.x[0:z + 2]) \
+ map (tofrom:s.y[3:3]) \
+ map (from: s.w, s.z[z + 1:z + 3], err)
{
err = 0;
- if (s.u != 1 || s.v.b.y != 3LL)
+ if (s.u != 1 || s.v.b.y != 3LL || s.x[0] != 7 || s.x[1] != 8
+ || s.y[3] != 9 || s.y[4] != 10 || s.y[5] != 11)
err = 1;
s.w.x = 6;
+ s.y[3] = 12;
+ s.y[4] = 13;
+ s.y[5] = 14;
+ s.z[1] = 15;
+ s.z[2] = 16;
+ s.z[3] = 17;
}
- if (err || s.w.x != 6)
+ if (err || s.w.x != 6 || s.y[3] != 12 || s.y[4] != 13 || s.y[5] != 14
+ || s.z[1] != 15 || s.z[2] != 16 || s.z[3] != 17)
abort ();
s.u++;
s.v.a++;
s.v.b.y++;
s.w.x++;
+ s.x[1] = 18;
+ s.z[0] = 19;
#pragma omp target data map (tofrom: s)
- #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ #pragma omp target map (always to: s.w, s.x[1], err) map (alloc:s.u, s.v.b,
s.z[z:z + 1])
{
err = 0;
- if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7)
+ if (s.u != 2 || s.v.b.y != 4LL || s.w.x != 7 || s.x[1] != 18 || s.z[0] !=
19)
err = 1;
s.w.x = 8;
+ s.x[1] = 20;
+ s.z[0] = 21;
}
- if (err || s.w.x != 8)
+ if (err || s.w.x != 8 || s.x[1] != 20 || s.z[0] != 21)
abort ();
s.u++;
s.v.a++;
s.v.b.y++;
s.w.x++;
- #pragma omp target data map (from: s.w) map (to: s.v.b, s.u)
- #pragma omp target map (always to: s.w, err) map (alloc:s.u, s.v.b)
+ s.x[0] = 22;
+ s.x[1] = 23;
+ #pragma omp target data map (from: s.w, s.x[0:2]) map (to: s.v.b, s.u)
+ #pragma omp target map (always to: s.w, s.x[0:2], err) map (alloc:s.u, s.v.b)
{
err = 0;
- if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9)
+ if (s.u != 3 || s.v.b.y != 5LL || s.w.x != 9 || s.x[0] != 22 || s.x[1] !=
23)
err = 1;
s.w.x = 11;
+ s.x[0] = 24;
+ s.x[1] = 25;
}
- if (err || s.w.x != 11)
+ if (err || s.w.x != 11 || s.x[0] != 24 || s.x[1] != 25)
abort ();
return 0;
}
--- libgomp/testsuite/libgomp.c/target-22.c.jj 2015-08-27 13:13:09.999364928
+0200
+++ libgomp/testsuite/libgomp.c/target-22.c 2015-08-28 19:58:50.109378650
+0200
@@ -0,0 +1,51 @@
+extern void abort (void);
+struct T { int a; int *b; int c; };
+struct S { int *s; char *u; struct T v; short *w; };
+volatile int z;
+
+int
+main ()
+{
+ struct S s;
+ int a[32], i;
+ char b[32];
+ short c[32];
+ for (i = 0; i < 32; i++)
+ {
+ a[i] = i;
+ b[i] = 32 + i;
+ c[i] = 64 + i;
+ }
+ s.s = a;
+ s.u = b + 2;
+ s.v.b = a + 16;
+ s.w = c + 3;
+ int err = 0;
+ #pragma omp target map (to:s.v.b[0:z + 7], s.u[z + 1:z + 4]) \
+ map (tofrom:s.s[3:3]) \
+ map (from: s.w[z:4], err) private (i)
+ {
+ err = 0;
+ for (i = 0; i < 7; i++)
+ if (s.v.b[i] != 16 + i)
+ err = 1;
+ for (i = 1; i < 5; i++)
+ if (s.u[i] != 34 + i)
+ err = 1;
+ for (i = 3; i < 6; i++)
+ if (s.s[i] != i)
+ err = 1;
+ else
+ s.s[i] = 128 + i;
+ for (i = 0; i < 4; i++)
+ s.w[i] = 96 + i;
+ }
+ if (err)
+ abort ();
+ for (i = 0; i < 32; i++)
+ if (a[i] != ((i >= 3 && i < 6) ? 128 + i : i)
+ || b[i] != 32 + i
+ || c[i] != ((i >= 3 && i < 7) ? 93 + i : 64 + i))
+ abort ();
+ return 0;
+}
--- libgomp/testsuite/libgomp.c/target-23.c.jj 2015-08-31 14:09:40.386455884
+0200
+++ libgomp/testsuite/libgomp.c/target-23.c 2015-08-31 14:10:33.475729499
+0200
@@ -0,0 +1,48 @@
+extern void abort (void);
+struct S { int s; int *u; int v[5]; };
+volatile int z;
+
+int
+main ()
+{
+ int u[10] = { 0, 1, 2, 3, 4, 5, 6, 7, 8, 9 }, err = 0;
+ struct S s = { 9, u + 3, { 10, 11, 12, 13, 14 } };
+ int *v = u + 4;
+ #pragma omp target enter data map (to: s.s, s.u[0:5]) map (alloc: s.v[1:3])
+ s.s++;
+ u[3]++;
+ s.v[1]++;
+ #pragma omp target update to (s.s) to (s.u[0:2], s.v[1:3])
+ #pragma omp target map (alloc: s.s, s.v[1:3]) map (from: err)
+ {
+ err = 0;
+ if (s.s != 10 || s.v[1] != 12 || s.v[2] != 12 || s.v[3] != 13)
+ err = 1;
+ if (v[-1] != 4 || v[0] != 4 || v[1] != 5 || v[2] != 6 || v[3] != 7)
+ err = 1;
+ s.s++;
+ s.v[2] += 2;
+ v[-1] = 5;
+ v[3] = 9;
+ }
+ if (err)
+ abort ();
+ #pragma omp target map (alloc: s.u[0:5])
+ {
+ err = 0;
+ if (s.u[0] != 5 || s.u[1] != 4 || s.u[2] != 5 || s.u[3] != 6 || s.u[4] !=
9)
+ err = 1;
+ s.u[1] = 12;
+ }
+ #pragma omp target update from (s.s, s.u[0:5]) from (s.v[1:3])
+ if (err || s.s != 11 || u[0] != 0 || u[1] != 1 || u[2] != 2 || u[3] != 5
+ || u[4] != 12 || u[5] != 5 || u[6] != 6 || u[7] != 9 || u[8] != 8
+ || u[9] != 9 || s.v[0] != 10 || s.v[1] != 12 || s.v[2] != 14
+ || s.v[3] != 13 || s.v[4] != 14)
+ abort ();
+ #pragma omp target exit data map (release: s.s)
+ #pragma omp target exit data map (release: s.u[0:5])
+ #pragma omp target exit data map (delete: s.v[1:3])
+ #pragma omp target exit data map (release: s.s)
+ return 0;
+}
Jakub
