@@ -18892,7 +18892,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)
@@ -18907,15 +18907,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++;
}
@@ -18923,6 +18943,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);
@@ -18964,10 +18985,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;
}
@@ -19016,8 +19051,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;
@@ -15811,6 +15811,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;
@@ -41719,16 +41719,34 @@ 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)
+ {
+ int n = cp_parser_skip_balanced_tokens (parser, pos + 1);
+ if (n != pos + 1)
+ {
+ 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++;
}
@@ -41736,6 +41754,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);
@@ -41785,10 +41804,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,
@@ -41852,8 +41890,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;
}
@@ -8375,6 +8375,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;
@@ -8906,7 +8906,12 @@ build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind)
gcc_assert (OMP_ITERATOR_DECL_P (t));
if (*last_bind)
- gimplify_and_add (*last_bind, pre_p);
+ {
+ 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;
+ }
tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
*last_bind = build3 (BIND_EXPR, void_type_node,
BLOCK_VARS (block), NULL, block);
@@ -9330,6 +9335,166 @@ 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;
+
+ /* Find all map iterators in use. */
+ hash_set<tree> all_iterators;
+ for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ all_iterators.add (TREE_PURPOSE (OMP_CLAUSE_DECL (c)));
+
+ /* Expand all clauses using the same iterator together. */
+ for (hash_set<tree>::iterator it = all_iterators.begin ();
+ it != all_iterators.end (); ++it)
+ {
+ for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c))
+ {
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP
+ || !OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ continue;
+
+ tree t = OMP_CLAUSE_DECL (c);
+ if (TREE_PURPOSE (t) != *it)
+ continue;
+
+ if (!check_iterator_var_usage (c))
+ continue;
+
+ if (TREE_PURPOSE (t) != last_iter)
+ {
+ tree tcnt = compute_iterator_count (t, pre_p);
+ if (!tcnt)
+ 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_int (SIZE_MAX);
+ }
+ }
+
+ 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
@@ -9538,6 +9703,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. */
@@ -9592,7 +9773,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");
@@ -9604,7 +9785,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;
@@ -9616,7 +9797,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;
@@ -9670,7 +9851,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)
@@ -9973,7 +10156,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
@@ -10160,7 +10343,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)
{
@@ -10699,7 +10882,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);
@@ -10886,7 +11069,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);
@@ -11208,7 +11393,7 @@ omp_accumulate_sibling_list (enum omp_region_type region_type,
using namespace omp_addr_tokenizer;
poly_offset_int coffset;
poly_int64 cbitpos;
- tree ocd = OMP_CLAUSE_DECL (grp_end);
+ tree ocd = OMP_ITERATOR_CLAUSE_DECL (grp_end);
bool openmp = !(region_type & ORT_ACC);
bool target = (region_type & ORT_TARGET) != 0;
tree *continue_at = NULL;
@@ -11806,7 +11991,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);
@@ -12091,6 +12276,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;
+ tree last_iterators = NULL_TREE;
int nowait = -1;
ctx = new_omp_context (region_type);
@@ -12471,6 +12657,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
grpnum++;
}
decl = OMP_CLAUSE_DECL (c);
+ last_iterators =
+ OMP_ITERATOR_DECL_P (decl) ? TREE_PURPOSE (decl) : NULL_TREE;
+ OMP_CLAUSE_DECL (c) = decl = OMP_ITERATOR_CLAUSE_DECL (c);
if (error_operand_p (decl))
{
@@ -13388,6 +13577,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
remove = true;
+ if (last_iterators)
+ OMP_CLAUSE_DECL (c) = build_tree_list (last_iterators,
+ OMP_CLAUSE_DECL (c));
if (remove)
*list_p = OMP_CLAUSE_CHAIN (c);
else
@@ -14106,7 +14298,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
break;
case OMP_CLAUSE_MAP:
- decl = OMP_CLAUSE_DECL (c);
+ decl = OMP_ITERATOR_CLAUSE_DECL (c);
if (!grp_end)
{
grp_start_p = list_p;
@@ -14168,7 +14360,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
: TYPE_SIZE_UNIT (TREE_TYPE (decl));
}
gimplify_omp_ctxp = ctx->outer_context;
- if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ /* Gimplify the OMP_CLAUSE_SIZE later, when the iterator is
+ gimplified. */
+ ;
+ else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR)
{
gimplify_omp_ctxp = ctx;
@@ -14333,6 +14529,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c))
break;
+ /* Do not gimplify the declaration yet for clauses with
+ iterators. */
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ break;
+
gimplify_omp_ctxp = ctx->outer_context;
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
fb_lvalue) == GS_ERROR)
@@ -14701,6 +14902,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
break;
}
+ gimplify_omp_map_iterators (orig_list_p, pre_p);
+
gimplify_omp_ctxp = ctx->outer_context;
delete_omp_context (ctx);
}
@@ -1518,7 +1518,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx)
case OMP_CLAUSE_MAP:
if (ctx->outer)
scan_omp_op (&OMP_CLAUSE_SIZE (c), ctx->outer);
- decl = OMP_CLAUSE_DECL (c);
+ decl = OMP_ITERATOR_CLAUSE_DECL (c);
/* If requested, make 'decl' addressable. */
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_DECL_MAKE_ADDRESSABLE (c))
@@ -12734,11 +12734,17 @@ 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)
+ {
+ 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,7 +13025,7 @@ 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)
@@ -13039,6 +13045,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
{
tree x = build_sender_ref (ovar, ctx);
tree v = ovar;
+ if (OMP_ITERATOR_DECL_P (OMP_CLAUSE_DECL (c)))
+ {
+ tree iterator = TREE_PURPOSE (OMP_CLAUSE_DECL (c));
+ v = TREE_VEC_ELT (iterator, 1);
+ }
if (in_reduction_clauses
&& OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& OMP_CLAUSE_MAP_IN_REDUCTION (c))
@@ -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,19 @@
+/* { 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" } */
+ /* Add a reference to x to ensure that the 'to' clause does not get
+ dropped. */
+ x[0] = 0;
+
+ #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))
{
@@ -2162,6 +2162,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;
@@ -1876,6 +1951,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 implicit clauses to map x). During gimplification, clauses with iterators are treated similarly to normal clauses, removing the iterator from the clause decl if necessary. gimplify_omp_map_iterators is called at the end of gimplify_adjust_omp_clauses. 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. The original OMP_CLAUSE_SIZE is set to SIZE_MAX to indicate that the clause has an expanded iterator associated with it. Multiple clauses using the same iterator are expanded together even if they are not adjacent. When OMP lowering clauses with iterators, the data array holding the expanded iterator info is used as the variable to send. 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 dd65c671dc9f5fb34290938a413c610eb0110ef6 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung <kcyeung@baylibre.com> Date: Mon, 2 Sep 2024 19:33:47 +0100 Subject: [PATCH 2/5] 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-09-02 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 (build_iterator_loop): Do not gimplify last binding into SSA. (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_accumulate_sibling_list): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. (omp_build_struct_sibling_lists): Likewise. (gimplify_scan_omp_clauses): Remove iterators from clauses before scanning clauses. Replace afterwards. (gimplify_adjust_omp_clauses): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. Skip gimplification of clause decl and size for clauses with iterators. Call gimplify_omp_map_iterators. * omp-low.cc (scan_sharing_clauses): Add field for iterator clauses. (lower_omp_target): Replace OMP_CLAUSE_DECL with OMP_ITERATOR_CLAUSE_DECL. Always increase map count by one for clauses with iterators. Use expanded iterator array as the output variable 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 | 63 ++++- gcc/cp/semantics.cc | 65 +++++ gcc/gimplify.cc | 229 +++++++++++++++++- gcc/omp-low.cc | 17 +- 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 | 19 ++ .../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, 793 insertions(+), 35 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