@@ -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,20 @@ 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);
+ OMP_CLAUSE_ITERATORS (c) = iterators;
+ }
parens.skip_until_found_close (parser);
return nl;
@@ -15058,7 +15058,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* We've reached the end of a list of expanded nodes. Reset the group
start pointer. */
if (c == grp_sentinel)
- grp_start_p = NULL;
+ {
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+ && OMP_CLAUSE_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc != grp_sentinel;
+ gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+ grp_start_p = NULL;
+ }
switch (OMP_CLAUSE_CODE (c))
{
@@ -15805,6 +15813,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_MAP:
if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
goto move_implicit;
+ if (OMP_CLAUSE_ITERATORS (c)
+ && c_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+ {
+ t = error_mark_node;
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -16497,6 +16511,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
pc = &OMP_CLAUSE_CHAIN (c);
}
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+ && OMP_CLAUSE_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
if (simdlen
&& safelen
&& tree_int_cst_lt (OMP_CLAUSE_SAFELEN_EXPR (safelen),
@@ -41776,16 +41776,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++;
}
@@ -41793,6 +41811,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);
@@ -41842,10 +41861,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,
@@ -41909,8 +41947,20 @@ 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);
+ OMP_CLAUSE_ITERATORS (c) = iterators;
+ }
return nlist;
}
@@ -7263,7 +7263,15 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
/* We've reached the end of a list of expanded nodes. Reset the group
start pointer. */
if (c == grp_sentinel)
- grp_start_p = NULL;
+ {
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+ && OMP_CLAUSE_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc != grp_sentinel;
+ gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+ grp_start_p = NULL;
+ }
switch (OMP_CLAUSE_CODE (c))
{
@@ -8484,6 +8492,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
case OMP_CLAUSE_MAP:
if (OMP_CLAUSE_MAP_IMPLICIT (c) && !implicit_moved)
goto move_implicit;
+ if (OMP_CLAUSE_ITERATORS (c)
+ && cp_omp_finish_iterators (OMP_CLAUSE_ITERATORS (c)))
+ {
+ t = error_mark_node;
+ break;
+ }
/* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
@@ -9348,6 +9362,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
pc = &OMP_CLAUSE_CHAIN (c);
}
+ if (grp_start_p
+ && OMP_CLAUSE_HAS_ITERATORS (*grp_start_p)
+ && OMP_CLAUSE_ITERATORS (*grp_start_p))
+ for (tree gc = *grp_start_p; gc; gc = OMP_CLAUSE_CHAIN (gc))
+ OMP_CLAUSE_ITERATORS (gc) = OMP_CLAUSE_ITERATORS (*grp_start_p);
+
if (reduction_seen < 0 && (ordered_seen || schedule_seen))
reduction_seen = -2;
@@ -8823,13 +8823,13 @@ gimplify_omp_affinity (tree *list_p, gimple_seq *pre_p)
}
/* Returns a tree expression containing the total iteration count of the
- iterator clause decl T. */
+ iterator IT. */
-static tree
-compute_iterator_count (tree t, gimple_seq *pre_p)
+tree
+compute_iterator_count (tree it, gimple_seq *pre_p)
{
tree tcnt = size_one_node;
- for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+ for (; it; it = TREE_CHAIN (it))
{
if (gimplify_expr (&TREE_VEC_ELT (it, 1), pre_p, NULL,
is_gimple_val, fb_rvalue) == GS_ERROR
@@ -8899,21 +8899,17 @@ compute_iterator_count (tree t, gimple_seq *pre_p)
Returns a pointer to the BIND_EXPR_BODY in the innermost loop body.
LAST_BIND is set to point to the BIND_EXPR containing the whole loop. */
-static tree *
-build_iterator_loop (tree c, gimple_seq *pre_p, tree *last_bind)
+tree *
+build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind)
{
- tree t = OMP_CLAUSE_DECL (c);
- gcc_assert (OMP_ITERATOR_DECL_P (t));
-
if (*last_bind)
gimplify_and_add (*last_bind, pre_p);
- tree block = TREE_VEC_ELT (TREE_PURPOSE (t), 5);
+ tree block = TREE_VEC_ELT (it, 5);
*last_bind = build3 (BIND_EXPR, void_type_node,
BLOCK_VARS (block), NULL, block);
TREE_SIDE_EFFECTS (*last_bind) = 1;
- SET_EXPR_LOCATION (*last_bind, OMP_CLAUSE_LOCATION (c));
tree *p = &BIND_EXPR_BODY (*last_bind);
- for (tree it = TREE_PURPOSE (t); it; it = TREE_CHAIN (it))
+ for (; it; it = TREE_CHAIN (it))
{
tree var = TREE_VEC_ELT (it, 0);
tree begin = TREE_VEC_ELT (it, 1);
@@ -9023,7 +9019,7 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
{
if (TREE_PURPOSE (t) != last_iter)
{
- tree tcnt = compute_iterator_count (t, pre_p);
+ tree tcnt = compute_iterator_count (TREE_PURPOSE (t), pre_p);
if (!tcnt)
return 2;
last_iter = TREE_PURPOSE (t);
@@ -9181,7 +9177,9 @@ gimplify_omp_depend (tree *list_p, gimple_seq *pre_p)
if (OMP_ITERATOR_DECL_P (t))
{
if (TREE_PURPOSE (t) != last_iter)
- last_body = build_iterator_loop (c, pre_p, &last_bind);
+ last_body = build_iterator_loop (TREE_PURPOSE (t), pre_p,
+ &last_bind);
+ SET_EXPR_LOCATION (last_bind, OMP_CLAUSE_LOCATION (c));
last_iter = TREE_PURPOSE (t);
if (TREE_CODE (TREE_VALUE (t)) == COMPOUND_EXPR)
{
@@ -12078,6 +12076,51 @@ error_out:
return success;
}
+/* Callback for walk_tree to find a VAR_DECL (stored in DATA) in the
+ tree TP. */
+
+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 that the clause C uses all the iterator variables.
+ Return TRUE if there are no errors. */
+
+static bool
+check_omp_map_iterators (tree c)
+{
+ bool error = false;
+ gcc_assert (OMP_CLAUSE_ITERATORS (c));
+
+ /* Do not check internal map kinds. */
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ return true;
+
+ for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it))
+ {
+ tree var = TREE_VEC_ELT (it, 0);
+ tree t = walk_tree (&OMP_CLAUSE_DECL (c), 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;
+}
+
/* Scan the OMP clauses in *LIST_P, installing mappings into a new
and previous omp contexts. */
@@ -12478,6 +12521,12 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
+ if (OMP_CLAUSE_ITERATORS (c) && !check_omp_map_iterators (c))
+ {
+ remove = true;
+ break;
+ }
+
if (!omp_parse_expr (addr_tokens, decl))
{
remove = true;
@@ -14168,7 +14217,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_CLAUSE_ITERATORS (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 +14386,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_CLAUSE_ITERATORS (c))
+ break;
+
gimplify_omp_ctxp = ctx->outer_context;
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue,
fb_lvalue) == GS_ERROR)
@@ -12607,6 +12607,163 @@ lower_omp_taskreg (gimple_stmt_iterator *gsi_p, omp_context *ctx)
}
}
+extern tree compute_iterator_count (tree it, gimple_seq *pre_p);
+extern tree *build_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind);
+
+struct iterator_loop_info_t
+{
+ tree bind;
+ tree count;
+ tree index;
+ tree *body;
+ tree *iterator;
+ hash_map<tree, tree> elems;
+};
+
+typedef hash_map<tree, iterator_loop_info_t> iterator_loop_map_t;
+
+/* Builds a loop to expand any iterators in clause C, reusing any previously
+ built loops if they use the same set of iterators. Generated Gimple
+ statements are placed into PRE_P. Information on the loops is held in
+ LOOPS. finish_omp_map_iterators must be called before the loops are
+ used. */
+
+static void
+lower_omp_map_iterators (tree c, gimple_seq *pre_p, iterator_loop_map_t *loops)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+ return;
+
+ bool built_p;
+ iterator_loop_info_t &loop = loops->get_or_insert (OMP_CLAUSE_ITERATORS (c),
+ &built_p);
+ if (!built_p)
+ {
+ loop.count = compute_iterator_count (OMP_CLAUSE_ITERATORS (c), pre_p);
+ if (!loop.count)
+ return;
+
+ loop.body = build_iterator_loop (OMP_CLAUSE_ITERATORS (c), pre_p,
+ &loop.bind);
+ loop.index = create_tmp_var (sizetype);
+ SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c));
+ loop.iterator = &OMP_CLAUSE_ITERATORS (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 (loop.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, loop.index,
+ size_binop (PLUS_EXPR, loop.index, size_int (2)));
+ append_to_statement_list_force (tem, loop.body);
+ }
+
+ /* Create array to hold expanded values. */
+ tree last_count_2 = size_binop (MULT_EXPR, loop.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");
+ }
+ loop.elems.put (c, elems);
+
+ /* 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, loop.count);
+ gimplify_and_add (tem, pre_p);
+}
+
+/* Set EXPR as the hostaddr expression that should result from the clause C.
+ LOOPS holds the intermediate loop info. Returns the tree that should be
+ passed as the hostaddr. */
+
+static tree
+lower_omp_map_iterator_expr (tree expr, tree c, iterator_loop_map_t *loops)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+ return expr;
+
+ iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c));
+ gcc_assert (loop);
+ tree *elems = loop->elems.get (c);
+
+ /* IN LOOP BODY: */
+ /* elems[idx] = <expr>; */
+ tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems, loop->index, NULL_TREE,
+ NULL_TREE);
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, void_type_node,
+ lhs, expr);
+ append_to_statement_list_force (tem, loop->body);
+
+ return build_fold_addr_expr_with_type (*elems, ptr_type_node);
+}
+
+/* Set SIZE as the size expression that should result from the clause C.
+ LOOPS holds the intermediate loop info. Returns the tree that should be
+ passed as the clause size. */
+
+static tree
+lower_omp_map_iterator_size (tree size, tree c, iterator_loop_map_t *loops)
+{
+ if (!OMP_CLAUSE_HAS_ITERATORS (c) || !OMP_CLAUSE_ITERATORS (c))
+ return size;
+
+ iterator_loop_info_t *loop = loops->get (OMP_CLAUSE_ITERATORS (c));
+ gcc_assert (loop);
+ tree *elems = loop->elems.get (c);
+
+ /* IN LOOP BODY: */
+ /* elems[idx+1] = size; */
+ tree lhs = build4 (ARRAY_REF, ptr_type_node, *elems,
+ size_binop (PLUS_EXPR, loop->index, size_int (1)),
+ NULL_TREE, NULL_TREE);
+ tree tem = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR,
+ void_type_node, lhs, size);
+ append_to_statement_list_force (tem, loop->body);
+
+ return size_int (SIZE_MAX);
+}
+
+/* Finish building the iterator loops in LOOPS, with generated Gimple
+ statements going in PRE_P. The loops cannot be amended after this is
+ called. */
+
+static void
+finish_omp_map_iterators (iterator_loop_map_t *loops, gimple_seq *pre_p)
+{
+ for (iterator_loop_map_t::iterator it = loops->begin ();
+ it != loops->end (); ++it)
+ {
+ iterator_loop_info_t &loop = (*it).second;
+ gimplify_and_add (loop.bind, pre_p);
+
+ for (hash_map<tree, tree>::iterator it2 = loop.elems.begin ();
+ it2 != loop.elems.end (); ++it2)
+ {
+ tree clause = (*it2).first;
+ OMP_CLAUSE_DECL (clause) = (*it2).second;
+ OMP_CLAUSE_SIZE (clause) = size_int (SIZE_MAX);
+ }
+ }
+}
+
/* Lower the GIMPLE_OMP_TARGET in the current statement
in GSI_P. CTX holds context information for the directive. */
@@ -12617,7 +12774,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tree child_fn, t, c;
gomp_target *stmt = as_a <gomp_target *> (gsi_stmt (*gsi_p));
gbind *tgt_bind, *bind, *dep_bind = NULL;
- gimple_seq tgt_body, olist, ilist, fplist, new_body;
+ gimple_seq tgt_body, olist, iterlist, ilist, fplist, new_body;
location_t loc = gimple_location (stmt);
bool offloaded, data_region;
unsigned int map_cnt = 0;
@@ -12628,6 +12785,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
tree deep_map_offset_data = NULL_TREE;
tree deep_map_offset = NULL_TREE;
+ iterator_loop_map_t iterator_loops;
+
offloaded = is_gimple_omp_offloaded (stmt);
switch (gimple_omp_target_kind (stmt))
{
@@ -12706,6 +12865,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
push_gimplify_context ();
fplist = NULL;
+ iterlist = NULL;
ilist = NULL;
olist = NULL;
for (c = clauses; c ; c = OMP_CLAUSE_CHAIN (c))
@@ -12761,7 +12921,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gcc_unreachable ();
}
#endif
- /* FALLTHRU */
+ lower_omp_map_iterators (c, &iterlist, &iterator_loops);
+ /* FALLTHRU */
case OMP_CLAUSE_TO:
case OMP_CLAUSE_FROM:
oacc_firstprivate:
@@ -13190,6 +13351,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
*p = build_fold_indirect_ref (nd);
}
v = build_fold_addr_expr_with_type (v, ptr_type_node);
+ v = lower_omp_map_iterator_expr (v, c, &iterator_loops);
gimplify_assign (x, v, &ilist);
nc = NULL_TREE;
}
@@ -13263,12 +13425,18 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
&& TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE)
{
gcc_assert (offloaded);
- tree avar
- = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
- mark_addressable (avar);
- gimplify_assign (avar, build_fold_addr_expr (var), &ilist);
- talign = DECL_ALIGN_UNIT (avar);
+ tree avar = build_fold_addr_expr (var);
+ if (!OMP_CLAUSE_ITERATORS (c))
+ {
+ tree tmp = create_tmp_var (TREE_TYPE (TREE_TYPE (x)));
+ mark_addressable (tmp);
+ gimplify_assign (tmp, avar, &ilist);
+ avar = tmp;
+ }
+ talign = TYPE_ALIGN_UNIT (TREE_TYPE (TREE_TYPE (x)));
avar = build_fold_addr_expr (avar);
+ avar = lower_omp_map_iterator_expr (avar, c,
+ &iterator_loops);
gimplify_assign (x, avar, &ilist);
}
else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FIRSTPRIVATE)
@@ -13348,6 +13516,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
if (s == NULL_TREE)
s = TYPE_SIZE_UNIT (TREE_TYPE (ovar));
s = fold_convert (size_type_node, s);
+ s = lower_omp_map_iterator_size (s, c, &iterator_loops);
purpose = size_int (map_idx++);
CONSTRUCTOR_APPEND_ELT (vsize, purpose, s);
if (TREE_CODE (s) != INTEGER_CST)
@@ -13713,6 +13882,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
DECL_INITIAL (TREE_VEC_ELT (t, 2))
= build_constructor (TREE_TYPE (TREE_VEC_ELT (t, 2)), vkind);
}
+ finish_omp_map_iterators (&iterator_loops, &iterlist);
for (int i = 1; i <= 2; i++)
if (deep_map_cnt || !TREE_STATIC (TREE_VEC_ELT (t, i)))
{
@@ -14280,6 +14450,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx)
gimple_omp_set_body (stmt, new_body);
}
+ gsi_insert_seq_before (gsi_p, iterlist, GSI_SAME_STMT);
+
bind = gimple_build_bind (NULL, NULL,
tgt_bind ? gimple_bind_block (tgt_bind)
: NULL_TREE);
@@ -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,23 @@
+/* { dg-do compile } */
+/* { dg-options "-fopenmp -fdump-tree-omplower" } */
+
+#define DIM1 10
+#define DIM2 20
+#define DIM3 30
+
+void f (int ***x, float ***y, double **z)
+{
+ #pragma omp target \
+ map(to: x, y) \
+ map(iterator(i=0:DIM1, j=0:DIM2), to: x[i][j][:DIM3], y[i][j][:DIM3]) \
+ map(from: z) \
+ 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 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "if \\(j <= 19\\) goto <D\.\[0-9\]+>; else goto <D\.\[0-9\]+>;" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):from:D\.\[0-9\]+" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1\\):attach:D\.\[0-9\]+" 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):to:D\.\[0-9\]+" 2 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1\\):attach:D\.\[0-9\]+" 4 "omplower" } } */
@@ -911,6 +911,11 @@ 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,");
+ if (OMP_CLAUSE_ITERATORS (clause))
+ {
+ dump_omp_iterators (pp, OMP_CLAUSE_ITERATORS (clause), spc, flags);
+ pp_colon (pp);
+ }
switch (OMP_CLAUSE_MAP_KIND (clause))
{
case GOMP_MAP_ALLOC:
@@ -266,7 +266,7 @@ unsigned const char omp_clause_num_ops[] =
1, /* OMP_CLAUSE_EXCLUSIVE */
2, /* OMP_CLAUSE_FROM */
2, /* OMP_CLAUSE_TO */
- 2, /* OMP_CLAUSE_MAP */
+ 3, /* OMP_CLAUSE_MAP */
1, /* OMP_CLAUSE_HAS_DEVICE_ADDR */
1, /* OMP_CLAUSE_DOACROSS */
2, /* OMP_CLAUSE__CACHE_ */
@@ -11598,6 +11598,9 @@ walk_tree_1 (tree *tp, walk_tree_fn func, void *data,
case OMP_CLAUSE:
{
int len = omp_clause_num_ops[OMP_CLAUSE_CODE (t)];
+ /* Do not walk the iterator operand of OpenMP MAP clauses. */
+ if (OMP_CLAUSE_HAS_ITERATORS (t))
+ len--;
for (int i = 0; i < len; i++)
WALK_SUBTREE (OMP_CLAUSE_OPERAND (t, i));
WALK_SUBTREE_TAIL (OMP_CLAUSE_CHAIN (t));
@@ -1619,6 +1619,13 @@ class auto_suppress_location_wrappers
!= UNKNOWN_LOCATION)
#define OMP_CLAUSE_LOCATION(NODE) (OMP_CLAUSE_CHECK (NODE))->omp_clause.locus
+#define OMP_CLAUSE_HAS_ITERATORS(NODE) \
+ (OMP_CLAUSE_CODE (NODE) == OMP_CLAUSE_MAP)
+#define OMP_CLAUSE_ITERATORS(NODE) \
+ OMP_CLAUSE_OPERAND (OMP_CLAUSE_RANGE_CHECK (OMP_CLAUSE_CHECK (NODE), \
+ OMP_CLAUSE_MAP, \
+ OMP_CLAUSE_MAP), 2)
+
/* True on OMP_FOR and other OpenMP/OpenACC looping constructs if the loop nest
is non-rectangular. */
#define OMP_FOR_NON_RECTANGULAR(NODE) \
@@ -975,6 +975,105 @@ gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
}
}
+static const char *
+kind_to_name (unsigned short kind)
+{
+ if (GOMP_MAP_IMPLICIT_P (kind))
+ kind &= ~GOMP_MAP_IMPLICIT;
+
+ switch (kind & 0xff)
+ {
+ case GOMP_MAP_ALLOC: return "GOMP_MAP_ALLOC";
+ case GOMP_MAP_FIRSTPRIVATE: return "GOMP_MAP_FIRSTPRIVATE";
+ case GOMP_MAP_FIRSTPRIVATE_INT: return "GOMP_MAP_FIRSTPRIVATE_INT";
+ case GOMP_MAP_TO: return "GOMP_MAP_TO";
+ case GOMP_MAP_TO_PSET: return "GOMP_MAP_TO_PSET";
+ case GOMP_MAP_FROM: return "GOMP_MAP_FROM";
+ case GOMP_MAP_TOFROM: return "GOMP_MAP_TOFROM";
+ case GOMP_MAP_POINTER: return "GOMP_MAP_POINTER";
+ case GOMP_MAP_ATTACH: return "GOMP_MAP_ATTACH";
+ case GOMP_MAP_DETACH: return "GOMP_MAP_DETACH";
+ default: return "unknown";
+ }
+}
+
+/* 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. ITERATOR_COUNT holds the iteration count of the
+ iterator that generates each map (0 if not generated from an iterator).
+ HOSTADDRS, SIZES, KINDS and ITERATOR_COUNT must be freed afterwards if any
+ merging occurs. */
+
+static bool
+gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes,
+ void **kinds, size_t **iterator_count)
+{
+ 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: %u -> %u\n",
+ (int) *mapnum, (int) 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;
+ *iterator_count = (size_t *) gomp_malloc (map_count * sizeof (size_t));
+
+ for (size_t i = 0; i < *mapnum; i++)
+ {
+ if ((*sizes)[i] == SIZE_MAX)
+ {
+ uintptr_t *iterator_array = (*hostaddrs)[i];
+ size_t count = *iterator_array++;
+ for (size_t j = 0; j < count; j++)
+ {
+ new_hostaddrs[new_idx] = (void *) *iterator_array++;
+ new_sizes[new_idx] = *iterator_array++;
+ new_kinds[new_idx] = (*skinds)[i];
+ (*iterator_count)[new_idx] = j + 1;
+ gomp_debug (1,
+ "Expanding map %u <%s>: "
+ "hostaddrs[%u] = %p, sizes[%u] = %lu\n",
+ (int) i, kind_to_name (new_kinds[new_idx]),
+ (int) new_idx, new_hostaddrs[new_idx],
+ (int) new_idx, (unsigned long) 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];
+ (*iterator_count)[new_idx] = 0;
+ 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,
@@ -991,6 +1090,11 @@ 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;
+ size_t *iterator_count = NULL;
+ if (short_mapkind)
+ iterators_p = gomp_merge_iterator_maps (&mapnum, &hostaddrs, &sizes,
+ &kinds, &iterator_count);
struct target_mem_desc *tgt
= gomp_malloc (sizeof (*tgt) + sizeof (tgt->list[0]) * mapnum);
tgt->list_count = mapnum;
@@ -1840,14 +1944,17 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
if (pragma_kind & GOMP_MAP_VARS_TARGET)
{
+ size_t map_num = 0;
for (i = 0; i < mapnum; i++)
- {
- cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
- gomp_copy_host2dev (devicep, aq,
- (void *) (tgt->tgt_start + i * sizeof (void *)),
- (void *) &cur_node.tgt_offset, sizeof (void *),
- true, cbufp);
- }
+ if (!iterator_count || iterator_count[i] <= 1)
+ {
+ cur_node.tgt_offset = gomp_map_val (tgt, hostaddrs, i);
+ gomp_copy_host2dev (devicep, aq,
+ (void *) (tgt->tgt_start + map_num * sizeof (void *)),
+ (void *) &cur_node.tgt_offset, sizeof (void *),
+ true, cbufp);
+ map_num++;
+ }
}
if (cbufp)
@@ -1879,6 +1986,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
}
gomp_mutex_unlock (&devicep->lock);
+
+ if (iterators_p)
+ {
+ free (hostaddrs);
+ free (sizes);
+ free (kinds);
+ free (iterator_count);
+ }
+
return tgt;
}
new file mode 100644
@@ -0,0 +1,47 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(to: 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,44 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(alloc: 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,56 @@
+/* { dg-do run } */
+/* { dg-require-effective-target offload_device_nonshared_as } */
+
+/* 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 enter data map(to: 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, storing it in the OMP_CLAUSE_ITERATOR argument of the clause. When finishing clauses, any clauses generated from a clause with iterators also has the iterator applied to them. During gimplification, check_omp_map_iterators is called to check that all iterator variables are referenced at some point with a clause. Gimplification of the clause decl and size are delayed until iterator expansion as they may reference iterator variables. In lower_target, lower_omp_map_iterators is called to construct the expansion loop for iterator clauses. Clauses using the same set of iterators reuse the loop, though with different storage allocated for them. lower_omp_map_iterator_expr is called to add the final expression that is sent as the hostaddr for libgomp to the loop, and a reference to the array generated by the iterator loop is returned to replace the original expression. lower_omp_map_iterator_size works similarly for the clause size. finish_omp_map_iterators is called later to finalise the loop. 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. As there are now multiple maps where one was previously, an entry is only added to the target vars for the first expanded map, otherwise it will get out of sync with the expected layout and the wrong variables will be picked up by the target function. From 50557e513ca534ba32f50d99991b056a07a6f671 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung <kcyeung@baylibre.com> Date: Fri, 4 Oct 2024 15:16:12 +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-10-04 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): Finish iterators. Apply iterators to generated clauses. gcc/cp/ * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. * semantics.cc (finish_omp_clauses): Finish iterators. Apply iterators to generated clauses. gcc/ * gimplify.cc (compute_iterator_count): Make non-static. Take an iterator instead of a clause for an operand. (build_iterator_loop): Likewise. (gimplify_omp_depend): Pass iterator in call to compute_iterator_count and build_iterator_loop. (find_var_decl): New. (check_omp_map_iterators): New. (gimplify_scan_omp_clauses): Call check_omp_map_iterators on clauses with iterators. (gimplify_adjust_omp_clauses): Skip gimplification of clause decl and size for clauses with iterators. * omp-low.cc (struct iterator_loop_info_t): New type. (iterator_loop_map_t): New type. (lower_omp_map_iterators): New. (lower_omp_map_iterator_expr): New. (lower_omp_map_iterator_size): New. (finish_omp_map_iterators): New. (lower_omp_target): Call lower_omp_map_iterators on clauses with iterators. Call lower_omp_map_iterator_expr before assigning to sender ref. Call lower_omp_map_iterator_size before setting the size. Call finish_omp_map_iterators. Insert statements generated during iterator expansion before the statements for the target clause. * tree-pretty-print.cc (dump_omp_clause): Call dump_omp_iterators for iterators in map clauses. * tree.cc (omp_clause_num_ops): Add operand for OMP_CLAUSE_MAP. (walk_tree_1): Do not walk last operand of OMP_CLAUSE_MAP. * tree.h (OMP_CLAUSE_HAS_ITERATORS): New. (OMP_CLAUSE_ITERATORS: New. gcc/testsuite/ * c-c++-common/gomp/map-6.c (foo): Amend expected error message. * c-c++-common/gomp/target-map-iterators-1.c: New. * c-c++-common/gomp/target-map-iterators-2.c: New. * c-c++-common/gomp/target-map-iterators-3.c: New. libgomp/ * target.c (kind_to_name): New. (gomp_merge_iterator_maps): New. (gomp_map_vars_internal): Call gomp_merge_iterator_maps. Copy address of only the first iteration to target vars. 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 | 59 +++++- gcc/c/c-typeck.cc | 22 ++- gcc/cp/parser.cc | 62 +++++- gcc/cp/semantics.cc | 22 ++- gcc/gimplify.cc | 88 +++++++-- gcc/omp-low.cc | 186 +++++++++++++++++- gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../gomp/target-map-iterators-1.c | 23 +++ .../gomp/target-map-iterators-2.c | 19 ++ .../gomp/target-map-iterators-3.c | 23 +++ gcc/tree-pretty-print.cc | 5 + gcc/tree.cc | 5 +- gcc/tree.h | 7 + libgomp/target.c | 130 +++++++++++- .../target-map-iterators-1.c | 47 +++++ .../target-map-iterators-2.c | 44 +++++ .../target-map-iterators-3.c | 56 ++++++ 17 files changed, 759 insertions(+), 49 deletions(-) create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/target-map-iterators-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