@@ -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));
@@ -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)
{
@@ -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)
@@ -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);
@@ -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;
@@ -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:
@@ -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 } */
+}
@@ -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;
+}
@@ -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);
}
@@ -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;
+}
@@ -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;
+}
@@ -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])
+}
@@ -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;
}
@@ -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;
+}
@@ -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;
+}