@@ -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,48 @@ 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 +6557,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 +6569,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 +6649,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),
@@ -7176,12 +7267,48 @@ 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))
- remove = true;
+ && ((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 (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_POINTER
@@ -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))
@@ -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;
+}
@@ -0,0 +1 @@
+#include "../libgomp.c/target-21.c"
@@ -0,0 +1,62 @@
+extern "C" void abort ();
+struct T { int a; int *b; int c; char (&d)[10]; };
+struct S { int *s; char *u; struct T v; short *w; short *&x; };
+volatile int z;
+
+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;
+ struct 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 ();
+ return 0;
+}
@@ -1,7 +1,12 @@
-extern void abort (void);
+extern
+#ifdef __cplusplus
+"C"
+#endif
+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 +18,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;
+}
Hi! 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. Jakub