https://gcc.gnu.org/g:28f09913235880d5011fb9b11a1c6e2c7ea6fcfb
commit 28f09913235880d5011fb9b11a1c6e2c7ea6fcfb Author: Kwok Cheung Yeung <kcye...@baylibre.com> Date: Wed Nov 27 21:49:32 2024 +0000 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. 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/ * gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded iterator loops. * gimple.cc (gimple_build_omp_target): Add argument for iterator loops sequence. Initialize iterator loops field. * gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET. * gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET. Add extra field for iterator loops. (gimple_build_omp_target): Add argument for iterator loops sequence. (gimple_omp_target_iterator_loops): New. (gimple_omp_target_iterator_loops_ptr): New. (gimple_omp_target_set_iterator_loops): New. * gimplify.cc (find_var_decl): New. (copy_omp_iterator): New. (remap_omp_iterator_var_1): New. (remap_omp_iterator_var): New. (remove_unused_omp_iterator_vars): New. (struct iterator_loop_info_t): New type. (iterator_loop_info_map_t): New type. (build_omp_iterators_loops): New. (enter_omp_iterator_loop_context_1): New. (enter_omp_iterator_loop_context): New. (enter_omp_iterator_loop_context): New. (exit_omp_iterator_loop_context): New. (gimplify_adjust_omp_clauses): Add argument for iterator loop sequence. Gimplify the clause decl and size into the iterator loop if iterators are used. (gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and build_omp_iterators_loops for OpenMP target expressions. Add loop sequence as argument when calling gimplify_adjust_omp_clauses and building the Gimple statement. * gimplify.h (enter_omp_iterator_loop_context): New prototype. (exit_omp_iterator_loop_context): New prototype. * gsstruct.def (GSS_OMP_TARGET): New. * omp-low.cc (lower_omp_map_iterator_expr): New. (lower_omp_map_iterator_size): New. (finish_omp_map_iterators): New. (lower_omp_target): Add sorry if iterators used with deep mapping. Call lower_omp_map_iterator_expr before assigning to sender ref. Call lower_omp_map_iterator_size before setting the size. Insert iterator loop sequence before the statements for the target clause. * tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator loop sequence of OpenMP target statements. (convert_local_reference_stmt): Likewise. (convert_tramp_reference_stmt): Likewise. * tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator information if present. (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. * c-c++-common/gomp/target-map-iterators-4.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. Co-authored-by: Andrew Stubbs <a...@baylibre.com> Diff: --- gcc/ChangeLog.omp | 55 +++ gcc/c/ChangeLog.omp | 6 + gcc/c/c-parser.cc | 56 ++- gcc/c/c-typeck.cc | 20 +- gcc/cp/ChangeLog.omp | 6 + gcc/cp/parser.cc | 57 ++- gcc/cp/semantics.cc | 20 +- gcc/gimple-pretty-print.cc | 6 + gcc/gimple.cc | 8 +- gcc/gimple.def | 2 +- gcc/gimple.h | 42 ++- gcc/gimplify.cc | 399 ++++++++++++++++++++- gcc/gimplify.h | 4 + gcc/gsstruct.def | 1 + gcc/omp-low.cc | 79 +++- gcc/testsuite/ChangeLog.omp | 8 + gcc/testsuite/c-c++-common/gomp/map-6.c | 10 +- .../c-c++-common/gomp/target-map-iterators-1.c | 23 ++ .../c-c++-common/gomp/target-map-iterators-2.c | 25 ++ .../c-c++-common/gomp/target-map-iterators-3.c | 23 ++ .../c-c++-common/gomp/target-map-iterators-4.c | 18 + gcc/tree-nested.cc | 8 + gcc/tree-pretty-print.cc | 14 + gcc/tree.cc | 5 +- gcc/tree.h | 8 + libgomp/ChangeLog.omp | 11 + libgomp/target.c | 130 ++++++- .../libgomp.c-c++-common/target-map-iterators-1.c | 47 +++ .../libgomp.c-c++-common/target-map-iterators-2.c | 44 +++ .../libgomp.c-c++-common/target-map-iterators-3.c | 56 +++ 30 files changed, 1141 insertions(+), 50 deletions(-) diff --git a/gcc/ChangeLog.omp b/gcc/ChangeLog.omp index cfdaa58a05d8..0d31c77933e0 100644 --- a/gcc/ChangeLog.omp +++ b/gcc/ChangeLog.omp @@ -1,3 +1,58 @@ +2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> + + * gimple-pretty-print.cc (dump_gimple_omp_target): Print expanded + iterator loops. + * gimple.cc (gimple_build_omp_target): Add argument for iterator + loops sequence. Initialize iterator loops field. + * gimple.def (GIMPLE_OMP_TARGET): Set GSS symbol to GSS_OMP_TARGET. + * gimple.h (gomp_target): Set GSS symbol to GSS_OMP_TARGET. Add extra + field for iterator loops. + (gimple_build_omp_target): Add argument for iterator loops sequence. + (gimple_omp_target_iterator_loops): New. + (gimple_omp_target_iterator_loops_ptr): New. + (gimple_omp_target_set_iterator_loops): New. + * gimplify.cc (find_var_decl): New. + (copy_omp_iterator): New. + (remap_omp_iterator_var_1): New. + (remap_omp_iterator_var): New. + (remove_unused_omp_iterator_vars): New. + (struct iterator_loop_info_t): New type. + (iterator_loop_info_map_t): New type. + (build_omp_iterators_loops): New. + (enter_omp_iterator_loop_context_1): New. + (enter_omp_iterator_loop_context): New. + (enter_omp_iterator_loop_context): New. + (exit_omp_iterator_loop_context): New. + (gimplify_adjust_omp_clauses): Add argument for iterator loop + sequence. Gimplify the clause decl and size into the iterator + loop if iterators are used. + (gimplify_omp_workshare): Call remove_unused_omp_iterator_vars and + build_omp_iterators_loops for OpenMP target expressions. Add + loop sequence as argument when calling gimplify_adjust_omp_clauses + and building the Gimple statement. + * gimplify.h (enter_omp_iterator_loop_context): New prototype. + (exit_omp_iterator_loop_context): New prototype. + * gsstruct.def (GSS_OMP_TARGET): New. + * omp-low.cc (lower_omp_map_iterator_expr): New. + (lower_omp_map_iterator_size): New. + (finish_omp_map_iterators): New. + (lower_omp_target): Add sorry if iterators used with deep mapping. + Call lower_omp_map_iterator_expr before assigning to sender ref. + Call lower_omp_map_iterator_size before setting the size. Insert + iterator loop sequence before the statements for the target clause. + * tree-nested.cc (convert_nonlocal_reference_stmt): Walk the iterator + loop sequence of OpenMP target statements. + (convert_local_reference_stmt): Likewise. + (convert_tramp_reference_stmt): Likewise. + * tree-pretty-print.cc (dump_omp_iterators): Dump extra iterator + information if present. + (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. + 2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> * gimplify.cc (gimplify_omp_affinity): Use OMP_ITERATOR_DECL_P. diff --git a/gcc/c/ChangeLog.omp b/gcc/c/ChangeLog.omp index 41e3815e069d..cc3a28c836db 100644 --- a/gcc/c/ChangeLog.omp +++ b/gcc/c/ChangeLog.omp @@ -1,3 +1,9 @@ +2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> + + * 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. + 2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> * c-typeck.cc (handle_omp_array_sections): Use OMP_ITERATOR_DECL_P. diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 4b84a14e6908..e0aafc405545 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -19443,7 +19443,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, enum gomp_map_kind kind) @@ -19457,15 +19457,35 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) 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++; else if (c_parser_peek_nth_token_raw (parser, pos + 1)->type == CPP_OPEN_PAREN) @@ -19487,6 +19507,7 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) int present_modifier = 0; int mapper_modifier = 0; tree mapper_name = NULL_TREE; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { c_token *tok = c_parser_peek_token (parser); @@ -19584,11 +19605,24 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) present_modifier++; c_parser_consume_token (parser); } + 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%>, %<mapper%> or " - "%<present%>"); + "than %<always%>, %<close%>, %<iterator%>, " + "%<mapper%> or %<present%>"); parens.skip_until_found_close (parser); return list; } @@ -19639,9 +19673,19 @@ c_parser_omp_clause_map (c_parser *parser, tree list, enum gomp_map_kind kind) tree last_new = NULL_TREE; + 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_ITERATORS (c) = iterators; last_new = c; } diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc index 4d4f73efa478..c75322e732e3 100644 --- a/gcc/c/c-typeck.cc +++ b/gcc/c/c-typeck.cc @@ -15303,7 +15303,14 @@ 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)) + 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)) { @@ -16185,6 +16192,12 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) break; + 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: @@ -16924,6 +16937,11 @@ 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)) + 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), diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index 09000189cc40..cbc4eb230486 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,9 @@ +2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> + + * parser.cc (cp_parser_omp_clause_map): Parse 'iterator' modifier. + * semantics.cc (finish_omp_clauses): Finish iterators. Apply + iterators to generated clauses. + 2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> * pt.cc (tsubst_omp_clause_decl): Use OMP_ITERATOR_DECL_P. diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index ffe056f0be27..1e9b7469b743 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -42606,16 +42606,34 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) 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++; else if (cp_lexer_peek_nth_token (parser->lexer, pos + 1)->type == CPP_OPEN_PAREN) @@ -42628,6 +42646,7 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) bool present_modifier = false; bool mapper_modifier = false; tree mapper_name = NULL_TREE; + tree iterators = NULL_TREE; for (int pos = 1; pos < map_kind_pos; ++pos) { cp_token *tok = cp_lexer_peek_token (parser->lexer); @@ -42744,11 +42763,29 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) present_modifier = true; cp_lexer_consume_token (parser->lexer); } + 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%>, " - "%<mapper%> or %<present%>"); + "%<iterator%>, %<mapper%> or %<present%>"); cp_parser_skip_to_closing_parenthesis (parser, /*recovering=*/true, /*or_comma=*/false, @@ -42814,9 +42851,19 @@ cp_parser_omp_clause_map (cp_parser *parser, tree list, enum gomp_map_kind kind) tree last_new = NULL_TREE; + 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_ITERATORS (c) = iterators; last_new = c; } diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 51cf8b960857..42a61574e9d2 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -7601,7 +7601,14 @@ 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)) + 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)) { @@ -8949,6 +8956,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_DIM || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_GRID_STRIDE) break; + 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: @@ -9861,6 +9874,11 @@ 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)) + 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; diff --git a/gcc/gimple-pretty-print.cc b/gcc/gimple-pretty-print.cc index 82461e837536..00e2b16f2673 100644 --- a/gcc/gimple-pretty-print.cc +++ b/gcc/gimple-pretty-print.cc @@ -1834,6 +1834,12 @@ dump_gimple_omp_target (pretty_printer *buffer, const gomp_target *gs, default: gcc_unreachable (); } + if (gimple_omp_target_iterator_loops (gs)) + { + pp_string (buffer, "// Expanded iterator loops for #pragma omp target\n"); + dump_gimple_seq (buffer, gimple_omp_target_iterator_loops (gs), spc, flags); + pp_newline (buffer); + } if (flags & TDF_RAW) { dump_gimple_fmt (buffer, spc, flags, "%G%s <%+BODY <%S>%nCLAUSES <", gs, diff --git a/gcc/gimple.cc b/gcc/gimple.cc index 300b50fb144a..54385994e53f 100644 --- a/gcc/gimple.cc +++ b/gcc/gimple.cc @@ -1268,10 +1268,13 @@ gimple_build_omp_interop (tree clauses) BODY is the sequence of statements that will be executed. KIND is the kind of the region. - CLAUSES are any of the construct's clauses. */ + CLAUSES are any of the construct's clauses. + ITERATOR_LOOPS is an optional sequence containing constructed loops + for OpenMP iterators. */ gomp_target * -gimple_build_omp_target (gimple_seq body, int kind, tree clauses) +gimple_build_omp_target (gimple_seq body, int kind, tree clauses, + gimple_seq iterator_loops) { gomp_target *p = as_a <gomp_target *> (gimple_alloc (GIMPLE_OMP_TARGET, 0)); @@ -1279,6 +1282,7 @@ gimple_build_omp_target (gimple_seq body, int kind, tree clauses) gimple_omp_set_body (p, body); gimple_omp_target_set_clauses (p, clauses); gimple_omp_target_set_kind (p, kind); + gimple_omp_target_set_iterator_loops (p, iterator_loops); return p; } diff --git a/gcc/gimple.def b/gcc/gimple.def index d0d6ab564e1d..fcd6041e13ac 100644 --- a/gcc/gimple.def +++ b/gcc/gimple.def @@ -393,7 +393,7 @@ DEFGSCODE(GIMPLE_OMP_SINGLE, "gimple_omp_single", GSS_OMP_SINGLE_LAYOUT) DATA_ARG is a vec of 3 local variables in the parent function containing data to be mapped to CHILD_FN. This is used to implement the MAP clauses. */ -DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_PARALLEL_LAYOUT) +DEFGSCODE(GIMPLE_OMP_TARGET, "gimple_omp_target", GSS_OMP_TARGET) /* GIMPLE_OMP_TEAMS <BODY, CLAUSES, CHILD_FN, DATA_ARG> represents #pragma omp teams diff --git a/gcc/gimple.h b/gcc/gimple.h index 0c8985a662aa..5c9c7f4f6f0b 100644 --- a/gcc/gimple.h +++ b/gcc/gimple.h @@ -682,11 +682,14 @@ struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) }; /* GIMPLE_OMP_TARGET */ -struct GTY((tag("GSS_OMP_PARALLEL_LAYOUT"))) +struct GTY((tag("GSS_OMP_TARGET"))) gomp_target : public gimple_statement_omp_parallel_layout { - /* No extra fields; adds invariant: - stmt->code == GIMPLE_OMP_TARGET. */ + /* [ WORD 1-10 ] : base class */ + + /* [ WORD 11 ] + Iterator loops. */ + gimple_seq iterator_loops; }; /* GIMPLE_OMP_TASK */ @@ -1667,7 +1670,7 @@ gomp_scan *gimple_build_omp_scan (gimple_seq, tree); gomp_sections *gimple_build_omp_sections (gimple_seq, tree); gimple *gimple_build_omp_sections_switch (void); gomp_single *gimple_build_omp_single (gimple_seq, tree); -gomp_target *gimple_build_omp_target (gimple_seq, int, tree); +gomp_target *gimple_build_omp_target (gimple_seq, int, tree, gimple_seq = NULL); gomp_teams *gimple_build_omp_teams (gimple_seq, tree); gomp_atomic_load *gimple_build_omp_atomic_load (tree, tree, enum omp_memory_order); @@ -6412,6 +6415,37 @@ gimple_omp_target_set_data_arg (gomp_target *omp_target_stmt, } +/* Return the Gimple sequence used to store loops for OpenMP iterators used + by OMP_TARGET_STMT. */ + +inline gimple_seq +gimple_omp_target_iterator_loops (const gomp_target *omp_target_stmt) +{ + return omp_target_stmt->iterator_loops; +} + + +/* Return a pointer to the Gimple sequence used to store loops for OpenMP + iterators used by OMP_TARGET_STMT. */ + +inline gimple_seq * +gimple_omp_target_iterator_loops_ptr (gomp_target *omp_target_stmt) +{ + return &omp_target_stmt->iterator_loops; +} + + +/* Set ITERATOR_LOOPS to be the Gimple sequence used to store loops + constructed for OpenMP iterators in OMP_TARGET_STMT. */ + +inline void +gimple_omp_target_set_iterator_loops (gomp_target *omp_target_stmt, + gimple_seq iterator_loops) +{ + omp_target_stmt->iterator_loops = iterator_loops; +} + + /* Return the clauses associated with OMP_TEAMS GS. */ inline tree diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index 6e02490f43a6..78678c14b93e 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -9680,6 +9680,367 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) return p; } + +/* 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) +{ + if (*tp == (tree) data) + return *tp; + + return NULL_TREE; +} + +/* Returns an element-by-element copy of OMP iterator tree IT. */ + +static tree +copy_omp_iterator (tree it, int elem_count = -1) +{ + if (elem_count < 0) + elem_count = TREE_VEC_LENGTH (it); + tree new_it = make_tree_vec (elem_count); + for (int i = 0; i < TREE_VEC_LENGTH (it); i++) + TREE_VEC_ELT (new_it, i) = TREE_VEC_ELT (it, i); + + return new_it; +} + +/* Helper function for walk_tree in remap_omp_iterator_var. */ + +static tree +remap_omp_iterator_var_1 (tree *tp, int *, void *data) +{ + tree old_var = ((tree *) data)[0]; + tree new_var = ((tree *) data)[1]; + + if (*tp == old_var) + *tp = new_var; + return NULL_TREE; +} + +/* Replace instances of OLD_VAR in TP with NEW_VAR. */ + +static void +remap_omp_iterator_var (tree *tp, tree old_var, tree new_var) +{ + tree vars[2] = { old_var, new_var }; + walk_tree (tp, remap_omp_iterator_var_1, vars, NULL); +} + +/* Scan through all clauses using OpenMP iterators in LIST_P. If any + clauses have iterators with variables that are not used by the clause + decl or size, issue a warning and replace the iterator with a copy with + the unused variables removed. */ + +static void +remove_unused_omp_iterator_vars (tree *list_p) +{ + auto_vec< vec<tree> > iter_vars; + auto_vec<tree> new_iterators; + + for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + continue; + auto_vec<tree> vars; + bool need_new_iterators = false; + 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) + { + need_new_iterators = true; + if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP + || OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH) + warning_at (OMP_CLAUSE_LOCATION (c), 0, + "iterator variable %qE not used in clause " + "expression", DECL_NAME (var)); + } + else + vars.safe_push (var); + } + if (!need_new_iterators) + continue; + if (need_new_iterators && vars.is_empty ()) + { + /* No iteration variables are used in the clause - remove the + iterator from the clause. */ + OMP_CLAUSE_ITERATORS (c) = NULL_TREE; + continue; + } + + /* If a new iterator has been created for the current set of used + iterator variables, then use that as the iterator. Otherwise, + create a new iterator for the current iterator variable set. */ + unsigned i; + for (i = 0; i < iter_vars.length (); i++) + { + if (vars.length () != iter_vars[i].length ()) + continue; + bool identical_p = true; + for (unsigned j = 0; j < vars.length () && identical_p; j++) + identical_p = vars[j] == iter_vars[i][j]; + + if (identical_p) + break; + } + if (i < iter_vars.length ()) + OMP_CLAUSE_ITERATORS (c) = new_iterators[i]; + else + { + tree new_iters = NULL_TREE; + tree *new_iters_p = &new_iters; + tree new_vars = NULL_TREE; + tree *new_vars_p = &new_vars; + i = 0; + for (tree it = OMP_CLAUSE_ITERATORS (c); it && i < vars.length(); + it = TREE_CHAIN (it)) + { + tree var = TREE_VEC_ELT (it, 0); + if (var == vars[i]) + { + *new_iters_p = copy_omp_iterator (it); + *new_vars_p = build_decl (OMP_CLAUSE_LOCATION (c), VAR_DECL, + DECL_NAME (var), TREE_TYPE (var)); + DECL_ARTIFICIAL (*new_vars_p) = 1; + DECL_CONTEXT (*new_vars_p) = DECL_CONTEXT (var); + TREE_VEC_ELT (*new_iters_p, 0) = *new_vars_p; + new_iters_p = &TREE_CHAIN (*new_iters_p); + new_vars_p = &DECL_CHAIN (*new_vars_p); + i++; + } + } + tree new_block = make_node (BLOCK); + BLOCK_VARS (new_block) = new_vars; + TREE_VEC_ELT (new_iters, 5) = new_block; + new_iterators.safe_push (new_iters); + iter_vars.safe_push (vars.copy ()); + OMP_CLAUSE_ITERATORS (c) = new_iters; + } + + /* Remap clause to use the new variables. */ + i = 0; + for (tree it = OMP_CLAUSE_ITERATORS (c); it; it = TREE_CHAIN (it)) + { + tree old_var = vars[i++]; + tree new_var = TREE_VEC_ELT (it, 0); + remap_omp_iterator_var (&OMP_CLAUSE_DECL (c), old_var, new_var); + remap_omp_iterator_var (&OMP_CLAUSE_SIZE (c), old_var, new_var); + } + } + + for (unsigned i = 0; i < iter_vars.length (); i++) + iter_vars[i].release (); +} + +struct iterator_loop_info_t +{ + tree bind; + tree count; + tree index; + tree body_label; + auto_vec<tree> clauses; +}; + +typedef hash_map<tree, iterator_loop_info_t> iterator_loop_info_map_t; + +/* Builds a loop to expand any OpenMP iterators in the clauses in LIST_P, + reusing any previously built loops if they use the same set of iterators. + Generated Gimple statements are placed into LOOPS_SEQ_P. The clause + iterators are updated with information on how and where to insert code into + the loop body. */ + +static void +build_omp_iterators_loops (tree *list_p, gimple_seq *loops_seq_p) +{ + iterator_loop_info_map_t loops; + + for (tree c = *list_p; c; c = OMP_CLAUSE_CHAIN (c)) + { + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + continue; + + bool built_p; + iterator_loop_info_t &loop + = loops.get_or_insert (OMP_CLAUSE_ITERATORS (c), &built_p); + + if (!built_p) + { + loop.count = compute_omp_iterator_count (OMP_CLAUSE_ITERATORS (c), + loops_seq_p); + if (!loop.count) + continue; + + loop.bind = NULL_TREE; + tree *body = build_omp_iterator_loop (OMP_CLAUSE_ITERATORS (c), + loops_seq_p, &loop.bind); + + loop.index = create_tmp_var (sizetype); + SET_EXPR_LOCATION (loop.bind, OMP_CLAUSE_LOCATION (c)); + + /* BEFORE LOOP: */ + /* idx = -1; */ + /* This should be initialized to before the individual elements, + as idx is pre-incremented in the loop body. */ + gimple *assign = gimple_build_assign (loop.index, size_int (-1)); + gimple_seq_add_stmt (loops_seq_p, assign); + + /* IN LOOP BODY: */ + /* Create a label so we can find this point later. */ + loop.body_label = create_artificial_label (OMP_CLAUSE_LOCATION (c)); + tree tem = build1 (LABEL_EXPR, void_type_node, loop.body_label); + append_to_statement_list_force (tem, body); + + /* idx += 2; */ + 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, 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, "omp_iter_data"); + TREE_ADDRESSABLE (elems) = 1; + gimple_add_tmp_var (elems); + } + else + { + /* Handle dynamic sizes. */ + sorry ("dynamic iterator sizes not implemented yet"); + } + + /* BEFORE LOOP: */ + /* 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, loops_seq_p); + + /* Make a copy of the iterator with extra info at the end. */ + int elem_count = TREE_VEC_LENGTH (OMP_CLAUSE_ITERATORS (c)); + tree new_iterator = copy_omp_iterator (OMP_CLAUSE_ITERATORS (c), + elem_count + 3); + TREE_VEC_ELT (new_iterator, elem_count) = loop.body_label; + TREE_VEC_ELT (new_iterator, elem_count + 1) = elems; + TREE_VEC_ELT (new_iterator, elem_count + 2) = loop.index; + TREE_CHAIN (new_iterator) = TREE_CHAIN (OMP_CLAUSE_ITERATORS (c)); + OMP_CLAUSE_ITERATORS (c) = new_iterator; + + loop.clauses.safe_push (c); + } + + /* Now gimplify and add all the loops that were built. */ + for (hash_map<tree, iterator_loop_info_t>::iterator it = loops.begin (); + it != loops.end (); ++it) + gimplify_and_add ((*it).second.bind, loops_seq_p); +} + +/* Helper function for enter_omp_iterator_loop_context. */ + +static gimple_seq * +enter_omp_iterator_loop_context_1 (tree iterator, gimple_seq *loops_seq_p) +{ + /* Drill into the nested bind expressions to get to the loop body. */ + for (gimple_stmt_iterator gsi = gsi_start (*loops_seq_p); + !gsi_end_p (gsi); gsi_next (&gsi)) + { + gimple *stmt = gsi_stmt (gsi); + + switch (gimple_code (stmt)) + { + case GIMPLE_BIND: + { + gbind *bind_stmt = as_a<gbind *> (stmt); + gimple_push_bind_expr (bind_stmt); + gimple_seq *bind_body_p = gimple_bind_body_ptr (bind_stmt); + gimple_seq *seq = + enter_omp_iterator_loop_context_1 (iterator, bind_body_p); + if (seq) + return seq; + gimple_pop_bind_expr (); + } + break; + case GIMPLE_TRY: + { + gimple_seq *try_eval_p = gimple_try_eval_ptr (stmt); + gimple_seq *seq = + enter_omp_iterator_loop_context_1 (iterator, try_eval_p); + if (seq) + return seq; + } + break; + case GIMPLE_LABEL: + { + glabel *label_stmt = as_a<glabel *> (stmt); + tree label = gimple_label_label (label_stmt); + if (label == TREE_VEC_ELT (iterator, 6)) + return loops_seq_p; + } + break; + default: + break; + } + } + + return NULL; +} + +/* Enter the Gimplification context in LOOPS_SEQ_P for the iterator loop + associated with OpenMP clause C. Returns the gimple_seq for the loop body + if C has OpenMP iterators, or ALT_SEQ_P if not. */ + +static gimple_seq * +enter_omp_iterator_loop_context (tree c, gimple_seq *loops_seq_p, + gimple_seq *alt_seq_p) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + return alt_seq_p; + + push_gimplify_context (); + + gimple_seq *seq = enter_omp_iterator_loop_context_1 (OMP_CLAUSE_ITERATORS (c), + loops_seq_p); + gcc_assert (seq); + return seq; +} + +/* Enter the Gimplification context in STMT for the iterator loop associated + with OpenMP clause C. Returns the gimple_seq for the loop body if C has + OpenMP iterators, or ALT_SEQ_P if not. */ + +gimple_seq * +enter_omp_iterator_loop_context (tree c, gomp_target *stmt, + gimple_seq *alt_seq_p) +{ + gimple_seq *loops_seq_p = gimple_omp_target_iterator_loops_ptr (stmt); + return enter_omp_iterator_loop_context (c, loops_seq_p, alt_seq_p); +} + +/* Exit the Gimplification context for the OpenMP clause C. */ + +void +exit_omp_iterator_loop_context (tree c) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + return; + while (!gimplify_ctxp->bind_expr_stack.is_empty ()) + gimple_pop_bind_expr (); + pop_gimplify_context (NULL); +} + /* If *LIST_P contains any OpenMP depend clauses with iterators, lower all the depend clauses by populating corresponding depend array. Returns 0 if there are no such depend clauses, or @@ -15253,7 +15614,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data) static void gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, - enum tree_code code) + enum tree_code code, + gimple_seq *loops_seq_p = NULL) { struct gimplify_omp_ctx *ctx = gimplify_omp_ctxp; tree *prev_list_p = NULL, *orig_list_p = list_p; @@ -15629,6 +15991,8 @@ 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; + gimple_seq *seq_p; + seq_p = enter_omp_iterator_loop_context (c, loops_seq_p, pre_p); if (GOMP_MAP_NONCONTIG_ARRAY_P (OMP_CLAUSE_MAP_KIND (c))) { gcc_assert (OMP_CLAUSE_SIZE (c) @@ -15637,12 +16001,12 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, of the individual array dimensions, which gimplify_expr doesn't handle, so skip the call to gimplify_expr here. */ } - else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), pre_p, NULL, + else if (gimplify_expr (&OMP_CLAUSE_SIZE (c), seq_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { gimplify_omp_ctxp = ctx; remove = true; - break; + goto end_adjust_omp_map_clause; } else if ((OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -15653,7 +16017,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, && TREE_CODE (OMP_CLAUSE_SIZE (c)) != INTEGER_CST) { OMP_CLAUSE_SIZE (c) - = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), pre_p, NULL, + = get_initialized_tmp_var (OMP_CLAUSE_SIZE (c), seq_p, NULL, false); if ((ctx->region_type & ORT_TARGET) != 0) omp_add_variable (ctx, OMP_CLAUSE_SIZE (c), @@ -15694,7 +16058,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, && (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)) { remove = true; - break; + goto end_adjust_omp_map_clause; } /* If we have a DECL_VALUE_EXPR (e.g. this is a class member and/or a variable captured in a lambda closure), look through that now @@ -15710,7 +16074,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, decl = OMP_CLAUSE_DECL (c) = DECL_VALUE_EXPR (decl); if (TREE_CODE (decl) == TARGET_EXPR) { - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL, is_gimple_lvalue, fb_lvalue) == GS_ERROR) remove = true; } @@ -15720,7 +16084,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* The OMP_CLAUSE_DECL for GRID_DIM/GRID_STRIDE isn't necessarily an lvalue -- e.g. it might be a constant. So handle it specially here. */ - if (gimplify_expr (&OMP_CLAUSE_DECL (c), pre_p, NULL, + if (gimplify_expr (&OMP_CLAUSE_DECL (c), seq_p, NULL, is_gimple_val, fb_rvalue) == GS_ERROR) { gimplify_omp_ctxp = ctx; @@ -15811,7 +16175,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* If we have e.g. map(struct: *var), don't gimplify the argument since omp-low.cc wants to see the decl itself. */ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT) - break; + goto end_adjust_omp_map_clause; /* If we have a non-contiguous (strided/rectangular) update operation with a VIEW_CONVERT_EXPR, we need to be careful not @@ -15826,10 +16190,10 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* We've already partly gimplified this in gimplify_scan_omp_clauses. Don't do any more. */ if (code == OMP_TARGET && OMP_CLAUSE_MAP_IN_REDUCTION (c)) - break; + goto end_adjust_omp_map_clause; gimplify_omp_ctxp = ctx->outer_context; - if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, + if (gimplify_expr (pd, seq_p, NULL, is_gimple_lvalue, fb_lvalue | fb_mayfail) == GS_ERROR) { sorry_at (OMP_CLAUSE_LOCATION (c), @@ -15847,7 +16211,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, GOVD_FIRSTPRIVATE | GOVD_SEEN); gimplify_omp_ctxp = ctx; - break; + goto end_adjust_omp_map_clause; } if ((code == OMP_TARGET @@ -15993,6 +16357,8 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, == GOMP_MAP_TO_PSET))) prev_list_p = list_p; +end_adjust_omp_map_clause: + exit_omp_iterator_loop_context (c); break; case OMP_CLAUSE_TO: @@ -18779,6 +19145,13 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) gcc_unreachable (); } + gimple_seq iterator_loops_seq = NULL; + if (TREE_CODE (expr) == OMP_TARGET) + { + remove_unused_omp_iterator_vars (&OMP_CLAUSES (expr)); + build_omp_iterators_loops (&OMP_CLAUSES (expr), &iterator_loops_seq); + } + bool save_in_omp_construct = in_omp_construct; if ((ort & ORT_ACC) == 0) in_omp_construct = false; @@ -18913,7 +19286,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) else gimplify_and_add (OMP_BODY (expr), &body); gimplify_adjust_omp_clauses (pre_p, body, &OMP_CLAUSES (expr), - TREE_CODE (expr)); + TREE_CODE (expr), &iterator_loops_seq); in_omp_construct = save_in_omp_construct; switch (TREE_CODE (expr)) @@ -18956,7 +19329,7 @@ gimplify_omp_workshare (tree *expr_p, gimple_seq *pre_p) break; case OMP_TARGET: stmt = gimple_build_omp_target (body, GF_OMP_TARGET_KIND_REGION, - OMP_CLAUSES (expr)); + OMP_CLAUSES (expr), iterator_loops_seq); break; case OMP_TARGET_DATA: /* Put use_device_{ptr,addr} clauses last, as map clauses are supposed diff --git a/gcc/gimplify.h b/gcc/gimplify.h index 5f31818e5b82..56125e819545 100644 --- a/gcc/gimplify.h +++ b/gcc/gimplify.h @@ -79,6 +79,10 @@ extern enum gimplify_status gimplify_expr (tree *, gimple_seq *, gimple_seq *, extern tree omp_get_construct_context (void); int omp_has_novariants (void); +extern gimple_seq *enter_omp_iterator_loop_context (tree, gomp_target *, + gimple_seq * = NULL); +extern void exit_omp_iterator_loop_context (tree); + extern void gimplify_type_sizes (tree, gimple_seq *); extern void gimplify_one_sizepos (tree *, gimple_seq *); extern gbind *gimplify_body (tree, bool); diff --git a/gcc/gsstruct.def b/gcc/gsstruct.def index 7708dc35fbfb..052cdd1a01a6 100644 --- a/gcc/gsstruct.def +++ b/gcc/gsstruct.def @@ -44,6 +44,7 @@ DEFGSSTRUCT(GSS_OMP, gimple_statement_omp, false) DEFGSSTRUCT(GSS_OMP_CRITICAL, gomp_critical, false) DEFGSSTRUCT(GSS_OMP_FOR, gomp_for, false) DEFGSSTRUCT(GSS_OMP_PARALLEL_LAYOUT, gimple_statement_omp_parallel_layout, false) +DEFGSSTRUCT(GSS_OMP_TARGET, gomp_target, false) DEFGSSTRUCT(GSS_OMP_TASK, gomp_task, false) DEFGSSTRUCT(GSS_OMP_SECTIONS, gomp_sections, false) DEFGSSTRUCT(GSS_OMP_SINGLE_LAYOUT, gimple_statement_omp_single_layout, false) diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc index d9735ff7c833..e87ea4664078 100644 --- a/gcc/omp-low.cc +++ b/gcc/omp-low.cc @@ -13658,6 +13658,61 @@ convert_from_firstprivate_int (tree var, tree orig_type, bool is_ref, return fold_build1 (VIEW_CONVERT_EXPR, type, tmp); } + /* 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, gomp_target *stmt) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + return expr; + + tree iterator = OMP_CLAUSE_ITERATORS (c); + tree elems = TREE_VEC_ELT (iterator, 7); + tree index = TREE_VEC_ELT (iterator, 8); + gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt); + + /* IN LOOP BODY: */ + /* elems[idx] = <expr>; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, index, + NULL_TREE, NULL_TREE); + tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, expr); + gimplify_and_add (mod_expr, loop_body_p); + exit_omp_iterator_loop_context (c); + + 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, gomp_target *stmt) +{ + if (!OMP_CLAUSE_HAS_ITERATORS (c)) + return size; + + tree iterator = OMP_CLAUSE_ITERATORS (c); + tree elems = TREE_VEC_ELT (iterator, 7); + tree index = TREE_VEC_ELT (iterator, 8); + gimple_seq *loop_body_p = enter_omp_iterator_loop_context (c, stmt); + + /* IN LOOP BODY: */ + /* elems[idx+1] = <size>; */ + tree lhs = build4 (ARRAY_REF, ptr_type_node, elems, + size_binop (PLUS_EXPR, index, size_int (1)), + NULL_TREE, NULL_TREE); + tree mod_expr = build2_loc (OMP_CLAUSE_LOCATION (c), MODIFY_EXPR, + void_type_node, lhs, size); + gimplify_and_add (mod_expr, loop_body_p); + exit_omp_iterator_loop_context (c); + + return size_int (SIZE_MAX); +} + /* Lower the GIMPLE_OMP_TARGET in the current statement in GSI_P. CTX holds context information for the directive. */ @@ -13857,6 +13912,11 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) deep_map_cnt = extra; } + if (deep_map_cnt + && OMP_CLAUSE_HAS_ITERATORS (c)) + sorry ("iterators used together with deep mapping are not " + "supported yet"); + if (!DECL_P (var)) { if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP @@ -14633,6 +14693,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, stmt); gimplify_assign (x, v, &ilist); nc = NULL_TREE; } @@ -14706,11 +14767,17 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) && TREE_CODE (TREE_TYPE (ovar)) == ARRAY_TYPE && 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, stmt); gimplify_assign (x, avar, &ilist); } else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP @@ -14837,6 +14904,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, stmt); purpose = size_int (map_idx++); CONSTRUCTOR_APPEND_ELT (vsize, purpose, s); if (TREE_CODE (s) != INTEGER_CST) @@ -15810,6 +15878,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) gimple_omp_set_body (stmt, new_body); } + gsi_insert_seq_before (gsi_p, gimple_omp_target_iterator_loops (stmt), + GSI_SAME_STMT); + gimple_omp_target_set_iterator_loops (stmt, NULL); bind = gimple_build_bind (NULL, NULL, tgt_bind ? gimple_bind_block (tgt_bind) : NULL_TREE); diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index 3b25acf21d08..bb5456a70eef 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,11 @@ +2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> + + * 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. + * c-c++-common/gomp/target-map-iterators-4.c: New. + 2025-04-17 Thomas Schwinge <tschwi...@baylibre.com> Backported from trunk: diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c index 6084ca26730c..dc9a8c662dbd 100644 --- a/gcc/testsuite/c-c++-common/gomp/map-6.c +++ b/gcc/testsuite/c-c++-common/gomp/map-6.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', 'mapper' or 'present'" } */ + #pragma omp target map (a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator', 'mapper' or 'present'" } */ ; - #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'mapper' or 'present'" } */ + #pragma omp target map (close, a to: b) /* { dg-error "'map' clause with map-type modifier other than 'always', 'close', 'iterator', 'mapper' 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', 'mapper' 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', 'mapper' 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', 'mapper' 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', 'mapper' 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', 'mapper' 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', 'mapper' or 'present'" } */ ; diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c new file mode 100644 index 000000000000..70076bd0a16d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-1.c @@ -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) + ; + + #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" } */ + ; +} diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c new file mode 100644 index 000000000000..57ebb1057063 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-2.c @@ -0,0 +1,25 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +void f (int *x, float *y, double *z) +{ + #pragma omp target map(iterator(i=0:10), to: x) /* { dg-warning "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(i2=0:10, j2=0:20), from: x[i2]) /* { dg-warning "iterator variable .j2. not used in clause expression" } */ + ; + + #pragma omp target map(iterator(i3=0:10, j3=0:20, k3=0:30), to: x[i3+j3], y[j3+k3], z[k3+i3]) + /* { dg-warning "iterator variable .i3. not used in clause expression" "" { target *-*-* } .-1 } */ + /* { dg-warning "iterator variable .j3. not used in clause expression" "" { target *-*-* } .-2 } */ + /* { dg-warning "iterator variable .k3. not used in clause expression" "" { target *-*-* } .-3 } */ + ; +} + +/* { dg-final { scan-tree-dump-times "map\\\(to:x" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i2=0:10:1, loop_label=\[^\\\)\]+\\\):from:" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i3=0:10:1, int j3=0:20:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int j3=0:20:1, int k3=0:30:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i3=0:10:1, int k3=0:30:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c new file mode 100644 index 000000000000..62df42ffde10 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-3.c @@ -0,0 +1,23 @@ +/* { 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(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\]+>;" 3 "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, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):from:\\*D\\\.\[0-9\]+" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):to:\\*D\\\.\[0-9\]+" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(iterator\\(int i=0:10:1, int j=0:20:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):attach:\\*D\\\.\[0-9\]+" 4 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c new file mode 100644 index 000000000000..5dc5ad51bfb3 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/target-map-iterators-4.c @@ -0,0 +1,18 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ +/* { dg-additional-options "-std=c++98" { target c++ } } */ + +int bar (int, int); +void baz (int, int *); +#pragma omp declare target enter (baz) + +void +foo (int x, int *p) +{ + #pragma omp target map (iterator (i=0:4), to: p[bar (x, i)]) + baz (x, p); +} + +/* { dg-final { scan-tree-dump "firstprivate\\\(x\\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump-times "bar \\\(x, i\\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\\(iterator\\\(int i=0:4:1, loop_label=" 2 "gimple" } } */ diff --git a/gcc/tree-nested.cc b/gcc/tree-nested.cc index e466b6f63d25..edd6ed3d96f7 100644 --- a/gcc/tree-nested.cc +++ b/gcc/tree-nested.cc @@ -1783,6 +1783,8 @@ convert_nonlocal_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + walk_body (convert_nonlocal_reference_stmt, convert_nonlocal_reference_op, + info, gimple_omp_target_iterator_loops_ptr (as_a <gomp_target *> (stmt))); if (!is_gimple_omp_offloaded (stmt)) { save_suppress = info->suppress_expansion; @@ -2516,6 +2518,9 @@ convert_local_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, break; case GIMPLE_OMP_TARGET: + walk_body (convert_local_reference_stmt, convert_local_reference_op, info, + gimple_omp_target_iterator_loops_ptr (as_a <gomp_target *> (stmt))); + if (!is_gimple_omp_offloaded (stmt)) { save_suppress = info->suppress_expansion; @@ -2912,6 +2917,9 @@ convert_tramp_reference_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p, case GIMPLE_OMP_TASK: do_parallel: { + if (gimple_code (stmt) == GIMPLE_OMP_TARGET) + walk_body (convert_tramp_reference_stmt, convert_tramp_reference_op, + info, gimple_omp_target_iterator_loops_ptr (as_a <gomp_target *> (stmt))); tree save_local_var_chain = info->new_local_var_chain; walk_gimple_op (stmt, convert_tramp_reference_op, wi); info->new_local_var_chain = NULL; diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 4371e9ac31f9..5a0c4fc43fc7 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -448,6 +448,15 @@ dump_omp_iterators (pretty_printer *pp, tree iter, int spc, dump_flags_t flags) pp_colon (pp); dump_generic_node (pp, TREE_VEC_ELT (it, 3), spc, flags, false); } + if (TREE_VEC_LENGTH (iter) > 6) + { + pp_string (pp, ", loop_label="); + dump_generic_node (pp, TREE_VEC_ELT (iter, 6), spc, flags, false); + pp_string (pp, ", elems="); + dump_generic_node (pp, TREE_VEC_ELT (iter, 7), spc, flags, false); + pp_string (pp, ", index="); + dump_generic_node (pp, TREE_VEC_ELT (iter, 8), spc, flags, false); + } pp_right_paren (pp); } @@ -1024,6 +1033,11 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) pp_string (pp, "readonly,"); if (OMP_CLAUSE_MAP_POINTS_TO_READONLY (clause)) pp_string (pp, "pt_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: diff --git a/gcc/tree.cc b/gcc/tree.cc index 4541cdbf20ab..5de1c3aa3815 100644 --- a/gcc/tree.cc +++ b/gcc/tree.cc @@ -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 */ 3, /* OMP_CLAUSE__MAPPER_BINDING_ */ @@ -11626,6 +11626,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)); diff --git a/gcc/tree.h b/gcc/tree.h index 7d123fa32f87..5dec4c0a8c62 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1634,6 +1634,14 @@ 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 \ + && OMP_CLAUSE_ITERATORS (NODE)) +#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) \ diff --git a/libgomp/ChangeLog.omp b/libgomp/ChangeLog.omp index 56ddb4378fd3..c93ae3b4fd5e 100644 --- a/libgomp/ChangeLog.omp +++ b/libgomp/ChangeLog.omp @@ -1,3 +1,14 @@ +2025-04-17 Kwok Cheung Yeung <kcye...@baylibre.com> + + * 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. + 2025-04-17 Thomas Schwinge <tschwi...@baylibre.com> * testsuite/libgomp.oacc-c++/exceptions-bad_cast-3.C: Adjust. diff --git a/libgomp/target.c b/libgomp/target.c index bdba51e30037..c0532ba1a673 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -976,6 +976,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, @@ -993,6 +1092,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 + nca_data_row_num)); @@ -2012,14 +2116,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) @@ -2051,6 +2158,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; } diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c new file mode 100644 index 000000000000..b3d87f231df1 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-1.c @@ -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; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c new file mode 100644 index 000000000000..8569b55ab5b5 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-2.c @@ -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; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c new file mode 100644 index 000000000000..be30fa65d807 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-3.c @@ -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; +}