@@ -18772,7 +18772,7 @@ c_parser_omp_clause_doacross (c_parser *parser, tree list)
map ( [map-type-modifier[,] ...] map-kind: variable-list )
map-type-modifier:
- always | close */
+ always | close | present | iterator (iterators-definition) */
static tree
c_parser_omp_clause_map (c_parser *parser, tree list)
@@ -18787,15 +18787,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int pos = 1;
int map_kind_pos = 0;
- while (c_parser_peek_nth_token_raw (parser, pos)->type == CPP_NAME)
+ int iterator_length = 0;
+ for (;;)
{
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COLON)
+ c_token *tok = c_parser_peek_nth_token_raw (parser, pos);
+ if (tok->type != CPP_NAME)
+ break;
+
+ const char *p = IDENTIFIER_POINTER (tok->value);
+ c_token *next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ if (strcmp (p, "iterator") == 0 && next_tok->type == CPP_OPEN_PAREN)
+ {
+ unsigned n = pos + 2;
+ if (c_parser_check_balanced_raw_token_sequence (parser, &n)
+ && c_parser_peek_nth_token_raw (parser, n)->type
+ == CPP_CLOSE_PAREN)
+ {
+ iterator_length = n - pos + 1;
+ pos = n;
+ next_tok = c_parser_peek_nth_token_raw (parser, pos + 1);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
pos++;
}
@@ -18803,6 +18823,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
int always_modifier = 0;
int close_modifier = 0;
int present_modifier = 0;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
c_token *tok = c_parser_peek_token (parser);
@@ -18844,10 +18865,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
}
present_modifier++;
}
+ else if (strcmp ("iterator", p) == 0
+ && c_parser_peek_2nd_token (parser)->type == CPP_OPEN_PAREN)
+ {
+ if (iterators)
+ {
+ c_parser_error (parser, "too many %<iterator%> modifiers");
+ parens.skip_until_found_close (parser);
+ return list;
+ }
+ iterators = c_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ continue;
+ }
else
{
c_parser_error (parser, "%<map%> clause with map-type modifier other "
- "than %<always%>, %<close%> or %<present%>");
+ "than %<always%>, %<close%>, %<iterator%> "
+ "or %<present%>");
parens.skip_until_found_close (parser);
return list;
}
@@ -18896,8 +18931,21 @@ c_parser_omp_clause_map (c_parser *parser, tree list)
nl = c_parser_omp_variable_list (parser, clause_loc, OMP_CLAUSE_MAP, list,
true);
+ if (iterators)
+ {
+ tree block = pop_scope ();
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (iterators)
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+ }
parens.skip_until_found_close (parser);
return nl;
@@ -15526,6 +15526,74 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ tree iterators = TREE_PURPOSE (t);
+ if (c_omp_finish_iterators (iterators))
+ {
+ t = error_mark_node;
+ break;
+ }
+
+ /* Find the end of the group of clauses that use the same
+ iterator.*/
+ tree end_clause;
+ for (end_clause = c; end_clause;
+ end_clause = OMP_CLAUSE_CHAIN (end_clause))
+ {
+ tree nc = OMP_CLAUSE_CHAIN (end_clause);
+ /* Remove iterator temporarily. */
+ OMP_CLAUSE_DECL (end_clause) =
+ TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+ if (!nc
+ || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+ || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+ break;
+ }
+ tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+ /* Temporarily split off the group of clauses with the same
+ iterator. */
+ OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+ tree new_clauses = c_finish_omp_clauses (c, ort);
+
+ /* Replace the iterators and splice the new clauses in. */
+ tree *clause_p = &new_clauses;
+ while (*clause_p)
+ {
+ /* Skip unwanted clause types.
+ FIXME: Is this the right thing to do? */
+ bool skip = false;
+ if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ skip = false;
+ break;
+ default:
+ skip = true;
+ break;
+ }
+ if (skip)
+ *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+ else
+ {
+ OMP_CLAUSE_DECL (*clause_p)
+ = build_tree_list (iterators,
+ OMP_CLAUSE_DECL (*clause_p));
+
+ clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+ }
+ }
+ *clause_p = next_clause;
+ *pc = new_clauses;
+ pc = clause_p;
+ continue;
+ }
+
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
grp_start_p = pc;
@@ -41672,16 +41672,35 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
int pos = 1;
int map_kind_pos = 0;
- while (cp_lexer_peek_nth_token (parser->lexer, pos)->type == CPP_NAME
- || cp_lexer_peek_nth_token (parser->lexer, pos)->keyword == RID_DELETE)
+ int iterator_length = 0;
+ for (;;)
{
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COLON)
+ cp_token *tok = cp_lexer_peek_nth_token (parser->lexer, pos);
+ if (!(tok->type == CPP_NAME || tok->keyword == RID_DELETE))
+ break;
+
+ cp_token *next_tok = cp_lexer_peek_nth_token (parser->lexer, pos + 1);
+ if (tok->type == CPP_NAME
+ && strcmp (IDENTIFIER_POINTER (tok->u.value), "iterator") == 0
+ && next_tok->type == CPP_OPEN_PAREN)
+ {
+ size_t n = cp_parser_skip_balanced_tokens (parser, 2);
+ if (cp_lexer_peek_nth_token (parser->lexer, n - 1)->type
+ == CPP_CLOSE_PAREN)
+ {
+ iterator_length = n - pos;
+ pos = n - 1;
+ next_tok = cp_lexer_peek_nth_token (parser->lexer, n);
+ }
+ }
+
+ if (next_tok->type == CPP_COLON)
{
map_kind_pos = pos;
break;
}
- if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_COMMA)
+ if (next_tok->type == CPP_COMMA)
pos++;
pos++;
}
@@ -41689,6 +41708,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
bool always_modifier = false;
bool close_modifier = false;
bool present_modifier = false;
+ tree iterators = NULL_TREE;
for (int pos = 1; pos < map_kind_pos; ++pos)
{
cp_token *tok = cp_lexer_peek_token (parser->lexer);
@@ -41738,10 +41758,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
}
present_modifier = true;
}
+ else if (strcmp ("iterator", p) == 0
+ && cp_lexer_peek_nth_token (parser->lexer, 2)->type
+ == CPP_OPEN_PAREN)
+ {
+ if (iterators)
+ {
+ cp_parser_error (parser, "too many %<iterator%> modifiers");
+ cp_parser_skip_to_closing_parenthesis (parser,
+ /*recovering=*/true,
+ /*or_comma=*/false,
+ /*consume_paren=*/true);
+ return list;
+ }
+ begin_scope (sk_omp, NULL);
+ iterators = cp_parser_omp_iterators (parser);
+ pos += iterator_length - 1;
+ continue;
+ }
else
{
cp_parser_error (parser, "%<map%> clause with map-type modifier other"
- " than %<always%>, %<close%> or %<present%>");
+ " than %<always%>, %<close%>, %<iterator%>"
+ " or %<present%>");
cp_parser_skip_to_closing_parenthesis (parser,
/*recovering=*/true,
/*or_comma=*/false,
@@ -41805,8 +41844,21 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list)
NULL, true);
finish_scope ();
+ if (iterators)
+ {
+ tree block = poplevel (1, 1, 0);
+ if (iterators == error_mark_node)
+ iterators = NULL_TREE;
+ else
+ TREE_VEC_ELT (iterators, 5) = block;
+ }
+
for (c = nlist; c != list; c = OMP_CLAUSE_CHAIN (c))
- OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ {
+ OMP_CLAUSE_SET_MAP_KIND (c, kind);
+ if (iterators)
+ OMP_CLAUSE_DECL (c) = build_tree_list (iterators, OMP_CLAUSE_DECL (c));
+ }
return nlist;
}
@@ -8362,6 +8362,71 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
auto_vec<omp_addr_token *, 10> addr_tokens;
t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ tree iterators = TREE_PURPOSE (t);
+ if (cp_omp_finish_iterators (iterators))
+ {
+ t = error_mark_node;
+ break;
+ }
+
+ /* Find the end of the group of clauses that use the same
+ iterator.*/
+ tree end_clause;
+ for (end_clause = c; end_clause;
+ end_clause = OMP_CLAUSE_CHAIN (end_clause))
+ {
+ tree nc = OMP_CLAUSE_CHAIN (end_clause);
+ /* Remove iterator temporarily. */
+ OMP_CLAUSE_DECL (end_clause) =
+ TREE_VALUE (OMP_CLAUSE_DECL (end_clause));
+ if (!nc
+ || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc))
+ || TREE_PURPOSE (OMP_CLAUSE_DECL (nc)) != iterators)
+ break;
+ }
+ tree next_clause = OMP_CLAUSE_CHAIN (end_clause);
+
+ /* Temporarily split off the group of clauses with the same
+ iterator. */
+ OMP_CLAUSE_CHAIN (end_clause) = NULL_TREE;
+ tree new_clauses = finish_omp_clauses (c, ort);
+
+ /* Replace the iterators and splice the new clauses in. */
+ tree *clause_p = &new_clauses;
+ while (*clause_p)
+ {
+ OMP_CLAUSE_DECL (*clause_p)
+ = build_tree_list (iterators,
+ OMP_CLAUSE_DECL (*clause_p));
+ /* Skip unwanted clause types.
+ FIXME: Is this the right thing to do? */
+ bool skip = false;
+ if (OMP_CLAUSE_CODE (*clause_p) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (*clause_p))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ skip = false;
+ break;
+ default:
+ skip = true;
+ break;
+ }
+ if (skip)
+ *clause_p = OMP_CLAUSE_CHAIN (*clause_p);
+ else
+ clause_p = &OMP_CLAUSE_CHAIN (*clause_p);
+ }
+ *clause_p = next_clause;
+ *pc = new_clauses;
+ pc = clause_p;
+ continue;
+ }
+
if (TREE_CODE (t) == OMP_ARRAY_SECTION)
{
grp_start_p = pc;
@@ -9253,6 +9253,172 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
return 1;
}
+/* Callback for walk_tree to find a VAR_DECL for the given tree. */
+
+static tree
+find_var_decl (tree *tp, int *, void *data)
+{
+ tree t = *tp;
+
+ if (TREE_CODE (t) == VAR_DECL && t == (tree) data)
+ return t;
+
+ return NULL_TREE;
+}
+
+/* Check for clause decls in iterators that do not use all the iterator
+ variables. */
+
+static bool
+check_iterator_var_usage (tree c)
+{
+ tree decl = OMP_CLAUSE_DECL (c);
+ bool error = false;
+ gcc_assert (OMP_ITERATOR_DECL_P (decl));
+
+ for (tree it = TREE_PURPOSE (decl); it; it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ tree t = walk_tree (&TREE_VALUE (decl), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ t = walk_tree (&OMP_CLAUSE_SIZE (c), find_var_decl, var, NULL);
+ if (t == NULL_TREE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "iterator variable %qD not used in clause expression",
+ var);
+ error = true;
+ }
+ }
+ return !error;
+}
+
+static void
+gimplify_omp_map_iterators (tree *list_p, gimple_seq *pre_p)
+{
+ tree last_iter = NULL_TREE;
+ tree last_bind = NULL_TREE;
+ tree last_count = NULL_TREE;
+ tree last_index = NULL_TREE;
+ tree *last_body = NULL;
+
+ while (tree c = *list_p)
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ {
+ list_p = &OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ tree t = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (t))
+ {
+ if (!check_iterator_var_usage (c))
+ {
+ *list_p = OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+
+ if (TREE_PURPOSE (t) != last_iter)
+ {
+ tree tcnt = compute_iterator_count (t, pre_p);
+ if (!tcnt)
+ {
+ *list_p = OMP_CLAUSE_CHAIN (c);
+ continue;
+ }
+ last_iter = TREE_PURPOSE (t);
+ last_count = tcnt;
+
+ last_body = build_iterator_loop (c, pre_p, &last_bind);
+ last_index = create_tmp_var (sizetype);
+ SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
+
+ /* idx = -1; */
+ /* This should be initialized to before the individual elements,
+ as idx is pre-incremented in the loop body. */
+ gimple *g = gimple_build_assign (last_index, size_int (-1));
+ gimple_seq_add_stmt (pre_p, g);
+
+ /* IN LOOP BODY: */
+ /* idx += 2; */
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, last_index,
+ size_binop (PLUS_EXPR, last_index,
+ size_int (2)));
+ append_to_statement_list_force (tem, last_body);
+ }
+
+ /* Create array to hold expanded values. */
+ tree last_count_2 = size_binop (MULT_EXPR, last_count, size_int (2));
+ tree arr_length = size_binop (PLUS_EXPR, last_count_2, size_int (1));
+ tree elems = NULL_TREE;
+ if (TREE_CONSTANT (arr_length))
+ {
+ tree type = build_array_type (ptr_type_node,
+ build_index_type (arr_length));
+ elems = create_tmp_var_raw (type);
+ TREE_ADDRESSABLE (elems) = 1;
+ gimple_add_tmp_var (elems);
+ }
+ else
+ {
+ /* Handle dynamic sizes. */
+ sorry ("Dynamic iterator sizes not implemented yet.");
+ }
+
+ /* elems[0] = count; */
+ tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, size_int (0),
+ NULL_TREE, NULL_TREE);
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, last_count);
+ gimplify_and_add (tem, pre_p);
+
+ /* IN LOOP BODY: */
+ /* elems[idx] = &<expr>; */
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems, last_index, NULL_TREE,
+ NULL_TREE);
+ tree rhs = build1 (ADDR_EXPR, ptr_type_node, TREE_VALUE (t));
+ tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, rhs);
+ append_to_statement_list_force (tem, last_body);
+
+ /* elems[idx+1] = OMP_CLAUSE_SIZE (c); */
+ lhs = build4 (ARRAY_REF, ptr_type_node, elems,
+ size_binop (PLUS_EXPR, last_index, size_int (1)),
+ NULL_TREE, NULL_TREE);
+ tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, OMP_CLAUSE_SIZE (c));
+ append_to_statement_list_force (tem, last_body);
+
+ /* Replace iterator information. */
+ TREE_PURPOSE (t) = make_tree_vec (2);
+ TREE_VEC_ELT (TREE_PURPOSE (t), 0) = last_iter;
+ TREE_VEC_ELT (TREE_PURPOSE (t), 1) = elems;
+
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ }
+ else if (last_bind)
+ {
+ bool saved_into_ssa = gimplify_ctxp->into_ssa;
+ gimplify_ctxp->into_ssa = false;
+ gimplify_and_add (last_bind, pre_p);
+ gimplify_ctxp->into_ssa = saved_into_ssa;
+ last_bind = NULL_TREE;
+ }
+
+ list_p = &OMP_CLAUSE_CHAIN (c);
+ }
+
+ if (last_bind)
+ {
+ bool saved_into_ssa = gimplify_ctxp->into_ssa;
+ gimplify_ctxp->into_ssa = false;
+ gimplify_and_add (last_bind, pre_p);
+ gimplify_ctxp->into_ssa = saved_into_ssa;
+ }
+}
+
/* True if mapping node C maps, or unmaps, a (Fortran) array descriptor. */
static bool
@@ -9461,6 +9627,22 @@ omp_get_base_pointer (tree expr)
return NULL_TREE;
}
+/* Return the iterator for a mapping group, or NULL if there isn't one. */
+
+static tree
+omp_group_iterator (omp_mapping_group *grp)
+{
+ tree c = grp->grp_end;
+ if (!OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ return NULL_TREE;
+
+ tree iter = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+ if (TREE_VEC_LENGTH (iter) == 2)
+ iter = TREE_VEC_ELT (iter, 0);
+
+ return iter;
+}
+
/* An attach or detach operation depends directly on the address being
attached/detached. Return that address, or none if there are no
attachments/detachments. */
@@ -9515,7 +9697,7 @@ omp_get_attachment (omp_mapping_group *grp)
case GOMP_MAP_ATTACH_DETACH:
case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
case GOMP_MAP_DETACH:
- return OMP_CLAUSE_DECL (node);
+ return OMP_ITERATOR_CLAUSE_DECL (node);
default:
internal_error ("unexpected mapping node");
@@ -9527,7 +9709,7 @@ omp_get_attachment (omp_mapping_group *grp)
node = OMP_CLAUSE_CHAIN (node);
if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_ATTACH
|| OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DETACH)
- return OMP_CLAUSE_DECL (node);
+ return OMP_ITERATOR_CLAUSE_DECL (node);
else
internal_error ("unexpected mapping node");
return error_mark_node;
@@ -9539,7 +9721,7 @@ omp_get_attachment (omp_mapping_group *grp)
return OMP_CLAUSE_DECL (*grp->grp_start);
if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
- return OMP_CLAUSE_DECL (*grp->grp_start);
+ return OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
else
internal_error ("unexpected mapping node");
return error_mark_node;
@@ -9593,7 +9775,9 @@ omp_group_last (tree *start_p)
== GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION)
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH
|| OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_ALWAYS_POINTER
- || omp_map_clause_descriptor_p (nc)))
+ || omp_map_clause_descriptor_p (nc))
+ && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+ == OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (nc)))
{
tree nc2 = OMP_CLAUSE_CHAIN (nc);
if (OMP_CLAUSE_MAP_KIND (nc) == GOMP_MAP_DETACH)
@@ -9896,7 +10080,7 @@ omp_index_mapping_groups_1 (hash_map<tree_operand_hash_no_se,
node && j < chained;
node = OMP_CLAUSE_CHAIN (node), j++)
{
- tree decl = OMP_CLAUSE_DECL (node);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (node);
/* Sometimes we see zero-offset MEM_REF instead of INDIRECT_REF,
meaning node-hash lookups don't work. This is a workaround for
that, but ideally we should just create the INDIRECT_REF at
@@ -10083,7 +10267,7 @@ omp_tsort_mapping_groups_1 (omp_mapping_group ***outlist,
}
}
- tree decl = OMP_CLAUSE_DECL (*grp->grp_start);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (*grp->grp_start);
while (decl)
{
@@ -10622,7 +10806,7 @@ omp_resolve_clause_dependencies (enum tree_code code,
FOR_EACH_VEC_ELT (*groups, i, grp)
{
tree grp_end = grp->grp_end;
- tree decl = OMP_CLAUSE_DECL (grp_end);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (grp_end);
gcc_assert (OMP_CLAUSE_CODE (grp_end) == OMP_CLAUSE_MAP);
@@ -10809,7 +10993,9 @@ omp_resolve_clause_dependencies (enum tree_code code,
{
omp_mapping_group *struct_group;
if (omp_mapped_by_containing_struct (grpmap, decl, &struct_group)
- && *grp->grp_start == grp_end)
+ && *grp->grp_start == grp_end
+ && omp_group_iterator (grp)
+ == omp_group_iterator (struct_group))
{
omp_check_mapping_compatibility (OMP_CLAUSE_LOCATION (grp_end),
struct_group, grp);
@@ -11729,7 +11915,7 @@ omp_build_struct_sibling_lists (enum tree_code code,
FOR_EACH_VEC_ELT (*groups, i, grp)
{
tree c = grp->grp_end;
- tree decl = OMP_CLAUSE_DECL (c);
+ tree decl = OMP_ITERATOR_CLAUSE_DECL (c);
tree grp_end = grp->grp_end;
auto_vec<omp_addr_token *> addr_tokens;
tree sentinel = OMP_CLAUSE_CHAIN (grp_end);
@@ -12014,6 +12200,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
tree c;
tree *orig_list_p = list_p;
int handled_depend_iterators = -1;
+ bool handled_map_iterators = false;
int nowait = -1;
ctx = new_omp_context (region_type);
@@ -12395,6 +12582,18 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
decl = OMP_CLAUSE_DECL (c);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ if (!handled_map_iterators)
+ {
+ gimplify_omp_map_iterators (list_p, pre_p);
+ handled_map_iterators = true;
+ continue;
+ }
+ /* Skip declarations with iterators. */
+ break;
+ }
+
if (error_operand_p (decl))
{
remove = true;
@@ -14076,6 +14275,9 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
if (remove)
break;
+ /* No further gimplfication required for clauses with iterators. */
+ if (OMP_ITERATOR_DECL_P (decl))
+ break;
if (OMP_CLAUSE_SIZE (c) == NULL_TREE)
{
/* Sanity check: attach/detach map kinds use the size as a bias,
@@ -1516,6 +1516,22 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_MAP:
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ /* FIXME: Is this the right way to handle these? */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ break;
+ tree field
+ = build_decl (OMP_CLAUSE_LOCATION (c),
+ FIELD_DECL, NULL_TREE, ptr_type_node);
+ SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node));
+ insert_field_into_struct (ctx->record_type, field);
+ splay_tree_insert (ctx->field_map,
+ (splay_tree_key) OMP_ITERATOR_CLAUSE_DECL (c),
+ (splay_tree_value) field);
+ break;
+ }
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
decl = OMP_CLAUSE_DECL (c);
@@ -12734,11 +12750,21 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
#endif
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c))
+ && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ {
+ /* Skip firstprivate pointers/references.
+ FIXME: Is this the right thing to do? */
+ if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE)
+ map_cnt++;
+ continue;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
oacc_firstprivate:
- var = OMP_CLAUSE_DECL (c);
+ var = OMP_ITERATOR_CLAUSE_DECL (c);
if (!DECL_P (var))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
@@ -13019,12 +13045,31 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
case OMP_CLAUSE_FROM:
oacc_firstprivate_map:
nc = c;
- ovar = OMP_CLAUSE_DECL (c);
+ ovar = OMP_ITERATOR_CLAUSE_DECL (c);
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_FIRSTPRIVATE_REFERENCE)))
break;
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ tree x = build_sender_ref (OMP_ITERATOR_CLAUSE_DECL (c), ctx);
+ tree iterator = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+ tree array = TREE_VEC_ELT (iterator, 1);
+ tree array_addr = build1 (ADDR_EXPR, ptr_type_node, array);
+ gimplify_assign (x, array_addr, &ilist);
+ purpose = size_int (map_idx++);
+ CONSTRUCTOR_APPEND_ELT (vsize, purpose, size_int (SIZE_MAX));
+
+ unsigned HOST_WIDE_INT tkind = OMP_CLAUSE_MAP_KIND (c);
+ gcc_checking_assert (tkind
+ < (HOST_WIDE_INT_C (1U) << talign_shift));
+ gcc_checking_assert (tkind
+ <= tree_to_uhwi (TYPE_MAX_VALUE (tkind_type)));
+ CONSTRUCTOR_APPEND_ELT (vkind, purpose,
+ build_int_cstu (tkind_type, tkind));
+ break;
+ }
if (!DECL_P (ovar))
{
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
@@ -13904,6 +13949,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
default:
break;
case OMP_CLAUSE_MAP:
+ /* FIXME: Handle firstprivate mappings for iterators properly. */
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ break;
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE)
{
@@ -13,19 +13,19 @@ foo (void)
#pragma omp target map (to:a)
;
- #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target enter data map(b7) map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target exit data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
- #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close' or 'present'" } */
+ #pragma omp target data map(b7) map (close, a from: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator' or 'present'" } */
;
new file mode 100644
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+#define DIM1 17
+#define DIM2 39
+
+void f (int **x, int **y)
+{
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2], y[i][:DIM2])
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2] + 2) /* { dg-message "unsupported map expression" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), iterator(j=0:DIM2), to: x[i][j]) /* { dg-error "too many 'iterator' modifiers" } */
+ ;
+
+ #pragma omp target map(iterator(i=0:DIM1), to: (i % 2 == 0) ? x[i] : y[i]) /* { dg-message "unsupported map expression" } */
+ ;
+}
new file mode 100644
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp" } */
+
+void f (int *x, float *y, double *z)
+{
+ #pragma omp target map(iterator(i=0:10), to: x) /* { dg-error "iterator variable .i. not used in clause expression" }*/
+ ;
+
+ #pragma omp target map(iterator(i=0:10, j=0:20), to: x[i]) /* { dg-error "iterator variable .j. not used in clause expression" }*/
+ ;
+
+ #pragma omp target map(iterator(i=0:10, j=0:20, k=0:30), to: x[i], y[j], z[k])
+ /* { dg-error "iterator variable .i. not used in clause expression" "" { target *-*-* } .-1 } */
+ /* { dg-error "iterator variable .j. not used in clause expression" "" { target *-*-* } .-2 } */
+ /* { dg-error "iterator variable .k. not used in clause expression" "" { target *-*-* } .-3 } */
+ ;
+}
new file mode 100644
@@ -0,0 +1,20 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-gimple" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+ #pragma omp target map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+ map(iterator(i=0:DIM1), from: z[i][:DIM2])
+ ;
+}
+
+/* { dg-final { scan-tree-dump-times "if \\(i <= 9\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:from:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):iterator_array=D\.\[0-9\]+:attach:" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:to:" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):iterator_array=D\.\[0-9\]+:attach:" 4 "gimple" } } */
@@ -451,6 +451,19 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags)
pp_right_paren (pp);
}
+static void
+dump_omp_map_iterators (pretty_printer *pp, tree iter, int spc,
+ dump_flags_t flags)
+{
+ if (TREE_VEC_LENGTH (iter) == 6)
+ dump_omp_iterators (pp, iter, spc, flags);
+ else
+ {
+ dump_omp_iterators (pp, TREE_VEC_ELT (iter, 0), spc, flags);
+ pp_string (pp, ":iterator_array=");
+ dump_generic_node (pp, TREE_VEC_ELT (iter, 1), spc, flags, false);
+ }
+}
/* Dump OMP clause CLAUSE, without following OMP_CLAUSE_CHAIN.
@@ -461,6 +474,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
{
const char *name;
const char *modifier = NULL;
+ tree decl = NULL_TREE;
switch (OMP_CLAUSE_CODE (clause))
{
case OMP_CLAUSE_PRIVATE:
@@ -911,6 +925,13 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
pp_string (pp, "map(");
if (OMP_CLAUSE_MAP_READONLY (clause))
pp_string (pp, "readonly,");
+ decl = OMP_CLAUSE_DECL (clause);
+ if (OMP_ITERATOR_DECL_P (decl))
+ {
+ dump_omp_map_iterators (pp, TREE_PURPOSE (decl), spc, flags);
+ pp_colon (pp);
+ decl = TREE_VALUE (decl);
+ }
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_ALLOC:
@@ -1025,8 +1046,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags)
gcc_unreachable ();
}
pp_colon (pp);
- dump_generic_node (pp, OMP_CLAUSE_DECL (clause),
- spc, flags, false);
+ dump_generic_node (pp, decl, spc, flags, false);
print_clause_size:
if (OMP_CLAUSE_SIZE (clause))
{
@@ -2153,6 +2153,13 @@ class auto_suppress_location_wrappers
&& TREE_PURPOSE (NODE) \
&& TREE_CODE (TREE_PURPOSE (NODE)) == TREE_VEC)
+/* Return the iterator expression if NODE contains an iterator.
+ Return the clause decl if NODE does not. */
+#define OMP_ITERATOR_CLAUSE_DECL(NODE) \
+ (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (NODE)) \
+ ? TREE_VALUE (OMP_CLAUSE_DECL (NODE)) \
+ : OMP_CLAUSE_DECL (NODE))
+
/* In a BLOCK (scope) node:
Variables declared in the scope NODE. */
#define BLOCK_VARS(NODE) (BLOCK_CHECK (NODE)->block.vars)
@@ -972,6 +972,77 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
}
}
+/* Map entries containing expanded iterators will be flattened and merged into
+ HOSTADDRS, SIZES and KINDS, and MAPNUM updated. Returns true if there are
+ any iterators found. HOSTADDRS, SIZES and KINDS must be freed afterwards
+ if any merging occurs. */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+ void **kinds)
+{
+ bool iterator_p = false;
+ size_t map_count = 0;
+ unsigned short **skinds = (unsigned short **) kinds;
+
+ for (size_t i = 0; i < *mapnum; i++)
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ map_count += iterator_array[0];
+ iterator_p = true;
+ }
+ else
+ map_count++;
+
+ if (!iterator_p)
+ return false;
+
+ gomp_debug (1,
+ "Expanding iterator maps - number of map entries: %ld -> %ld\n",
+ *mapnum, map_count);
+ void **new_hostaddrs = (void **) gomp_malloc (map_count * sizeof (void *));
+ size_t *new_sizes = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+ unsigned short *new_kinds
+ = (unsigned short *) gomp_malloc (map_count * sizeof (unsigned short));
+ size_t new_idx = 0;
+
+ for (size_t i = 0; i < *mapnum; i++)
+ {
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ size_t count = iterator_array[0];
+ for (int j = 1; j < count * 2 + 1; j += 2)
+ {
+ new_hostaddrs[new_idx] = (void *) iterator_array[j];
+ new_sizes[new_idx] = iterator_array[j+1];
+ new_kinds[new_idx] = (*skinds)[i];
+ gomp_debug (1,
+ "Expanding map %ld: "
+ "hostaddrs[%ld] = %p, sizes[%ld] = %ld\n",
+ i, new_idx, new_hostaddrs[new_idx],
+ new_idx, new_sizes[new_idx]);
+ new_idx++;
+ }
+ }
+ else
+ {
+ new_hostaddrs[new_idx] = (*hostaddrs)[i];
+ new_sizes[new_idx] = (*sizes)[i];
+ new_kinds[new_idx] = (*skinds)[i];
+ new_idx++;
+ }
+ }
+
+ *mapnum = map_count;
+ *hostaddrs = new_hostaddrs;
+ *sizes = new_sizes;
+ *kinds = new_kinds;
+
+ return true;
+}
+
static inline __attribute__((always_inline)) struct target_mem_desc *
gomp_map_vars_internal (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, size_t mapnum,
@@ -988,6 +1059,10 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
const int typemask = short_mapkind ? 0xff : 0x7;
struct splay_tree_s *mem_map = &devicep->mem_map;
struct splay_tree_key_s cur_node;
+ bool iterators_p = false;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds);
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
@@ -1873,6 +1948,14 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ }
+
return tgt;
}
new file mode 100644
@@ -0,0 +1,44 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+int mkarray (int *x[])
+{
+ int expected = 0;
+
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j];
+ }
+ }
+
+ return expected;
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y;
+
+ int expected = mkarray (x);
+
+ #pragma omp target map(iterator(i=0:DIM1), to: x[i][:DIM2]) map(from: y)
+ {
+ y = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+ }
+
+ return y - expected;
+}
new file mode 100644
@@ -0,0 +1,42 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays from target using map
+ iterators. */
+
+#include <stdlib.h>
+
+#define DIM1 8
+#define DIM2 15
+
+void mkarray (int *x[])
+{
+ for (int i = 0; i < DIM1; i++)
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+}
+
+int main (void)
+{
+ int *x[DIM1];
+ int y, expected;
+
+ mkarray (x);
+
+ #pragma omp target map(iterator(i=0:DIM1), from: x[i][:DIM2]) \
+ map(from: expected)
+ {
+ expected = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = (i+1) * (j+1);
+ expected += x[i][j];
+ }
+ }
+
+ y = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ y += x[i][j];
+
+ return y - expected;
+}
new file mode 100644
@@ -0,0 +1,54 @@
+/* { dg-do run } */
+
+/* Test transfer of dynamically-allocated arrays to target using map
+ iterators, with multiple iterators and function calls in the iterator
+ expression. */
+
+#include <stdlib.h>
+
+#define DIM1 16
+#define DIM2 15
+
+int mkarrays (int *x[], int *y[])
+{
+ int expected = 0;
+
+ for (int i = 0; i < DIM1; i++)
+ {
+ x[i] = (int *) malloc (DIM2 * sizeof (int));
+ y[i] = (int *) malloc (sizeof (int));
+ *y[i] = rand ();
+ for (int j = 0; j < DIM2; j++)
+ {
+ x[i][j] = rand ();
+ expected += x[i][j] * *y[i];
+ }
+ }
+
+ return expected;
+}
+
+int f (int i, int j)
+{
+ return i * 4 + j;
+}
+
+int main (void)
+{
+ int *x[DIM1], *y[DIM1];
+ int sum;
+
+ int expected = mkarrays (x, y);
+
+ #pragma omp target map(iterator(i=0:DIM1/4, j=0:4), to: x[f(i, j)][:DIM2]) \
+ map(iterator(i=0:DIM1), to: y[i][:1]) \
+ map(from: sum)
+ {
+ sum = 0;
+ for (int i = 0; i < DIM1; i++)
+ for (int j = 0; j < DIM2; j++)
+ sum += x[i][j] * y[i][0];
+ }
+
+ return sum - expected;
+}
This patch modifies the C and C++ parsers to accept an iterator as a map type modifier, encoded in the same way as the depend and affinity clauses. When finishing the clauses, clauses with iterators are treated separately from ones without to avoid clashes (e.g. iterating over x[i] will likely generate clauses to map x). During gimplification, gimplify_omp_map_iterators is called during scanning if a map clause encountered has any iterators. This scans all the remaining clauses in one go, as iterators may be shared between clauses. Later clauses with iterators are simply skipped over. For each map clause with an iterator, gimplify_omp_map_iterators generates a loop (or multiple loops, if the iterator is multidimensional) to iterate over the iterator expression, storing the result in a new array (constant-sized for now, we could dynamically allocate the array for non-constant iteration bounds). The data array stores the total number of iterations in the first element, then the address generated by the iterator expression and the OMP_CLAUSE_SIZE (since the iteration variables may occur within the size tree) for each iteration. The clause is then rewritten to point to the new array. The original clause decl is no longer directly relevant, but is kept around for informational purposes and to help with clause sorting. The original OMP_CLAUSE_SIZE is set to NULL_TREE. When OMP lowering clauses with iterators, the data array holding the expanded iterator info is allocated to a field in the omp_data, and the size is set to SIZE_MAX to mark the entry as coming from an expanded iterator. Libgomp has a new function gomp_merge_iterator_maps which identifies data coming from an iterator, and effectively creates new maps on-the-fly from the iterator info array, inserting them into the list of mappings at the point where iterator data occurred. From b2e8ff46929d5a2781781486ec942b344056d78b Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung <kcyeung@baylibre.com> Date: Tue, 12 Mar 2024 22:51:06 +0000 Subject: [PATCH 2/3] openmp: Add support for iterators in map clauses (C/C++) This adds preliminary support for iterators in map clauses within OpenMP 'target' constructs (which includes constructs such as 'target enter data'). Iterators with non-constant loop bounds are not currently supported. 2024-05-24 Kwok Cheung Yeung <kcyeung@baylibre.com> gcc/c/ * c-parser.cc (c_parser_omp_clause_map): Parse 'iterator' modifier. * c-typeck.cc (c_finish_omp_clauses): Call recursively on iterator clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Call recursively on iterator clauses. gcc/ * gimplify.cc (find_var_decl): New. (check_iterator_var_usage): New. (gimplify_omp_map_iterators): New. (omp_group_iterator): New. (omp_get_attachment): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_group_last): Keep decls with and without iterators in separate groups. (omp_index_mapping_groups_1): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_tsort_mapping_groups_1): Likewise. (omp_resolve_clause_dependencies): Likewise. Prevent removal of mapping if groups do not use the same iterators. (omp_build_struct_sibling_lists): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (gimplify_scan_omp_clauses): Call gimplify_omp_map_iterators once to handle clauses with iterators, then skip subsequent iterator clauses. * omp-low.cc (scan_sharing_clauses): Add field for iterator clauses. (lower_omp_target): Add map entries for iterator clauses. * tree-pretty-print.cc (dump_omp_map_iterators): New. (dump_omp_clause): Call dump_omp_map_iterators for iterators in map clauses. * tree.h (OMP_ITERATOR_CLAUSE_DECL): New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-iterator-1.c: New. * c-c++-common/gomp/target-iterator-2.c: New. * c-c++-common/gomp/target-iterator-3.c: New. libgomp/ * target.c (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Free allocated variables. * testsuite/libgomp.c-c++-common/target-map-iterators-1.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-2.c: New. * testsuite/libgomp.c-c++-common/target-map-iterators-3.c: New. --- gcc/c/c-parser.cc | 60 ++++- gcc/c/c-typeck.cc | 68 ++++++ gcc/cp/parser.cc | 64 ++++- gcc/cp/semantics.cc | 65 ++++++ gcc/gimplify.cc | 220 +++++++++++++++++- gcc/omp-low.cc | 52 ++++- gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../c-c++-common/gomp/target-iterator-1.c | 23 ++ .../c-c++-common/gomp/target-iterator-2.c | 17 ++ .../c-c++-common/gomp/target-iterator-3.c | 20 ++ gcc/tree-pretty-print.cc | 24 +- gcc/tree.h | 7 + libgomp/target.c | 83 +++++++ .../target-map-iterators-1.c | 44 ++++ .../target-map-iterators-2.c | 42 ++++ .../target-map-iterators-3.c | 54 +++++ 16 files changed, 823 insertions(+), 30 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-iterator-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c