https://gcc.gnu.org/g:17ae5b3d96a1a8a14922ad89c7f0c98f556d4b8a
commit 17ae5b3d96a1a8a14922ad89c7f0c98f556d4b8a Author: Julian Brown <[email protected]> Date: Wed Apr 29 12:12:13 2026 +0200 OpenMP: Expand "declare mapper" mappers for target {enter,exit,} data directives This patch allows 'declare mapper' mappers to be used on 'omp target data', 'omp target enter data' and 'omp target exit data' directives. For each of these, only explicit mappings are supported, unlike for 'omp target' directives where implicit uses of variables inside an offload region might trigger mappers also. Add support for C and C++. The patch also adjusts 'map kind decay' to match OpenMP 5.2 semantics, which is particularly important with regard to 'exit data' operations. gcc/c-family/ * c-common.h (c_omp_region_type): Add C_ORT_EXIT_DATA, C_ORT_OMP_EXIT_DATA. (c_omp_instantiate_mappers): Add region type parameter. * c-omp.cc (omp_split_map_kind, omp_join_map_kind, omp_map_decayed_kind): New functions. (omp_instantiate_mapper): Add ORT parameter. Implement map kind decay for instantiated mapper clauses. (c_omp_instantiate_mappers): Add ORT parameter, pass to omp_instantiate_mapper. gcc/c/ * c-parser.cc (c_parser_omp_target_data): Instantiate mappers for 'omp target data'. (c_parser_omp_target_enter_data): Instantiate mappers for 'omp target enter data'. (c_parser_omp_target_exit_data): Instantiate mappers for 'omp target exit data'. (c_parser_omp_target): Add c_omp_region_type argument to c_omp_instantiate_mappers call. * c-tree.h (c_omp_instantiate_mappers): Remove spurious prototype. gcc/cp/ * parser.cc (cp_parser_omp_target_data): Instantiate mappers for 'omp target data'. (cp_parser_omp_target_enter_data): Instantiate mappers for 'omp target enter data'. (cp_parser_omp_target_exit_data): Instantiate mappers for 'omp target exit data'. (cp_parser_omp_target): Add c_omp_region_type argument to c_omp_instantiate_mappers call. * pt.cc (tsubst_omp_clauses): Instantiate mappers for OMP regions other than just C_ORT_OMP_TARGET. (tsubst_expr): Update call to tsubst_omp_clauses for OMP_TARGET_UPDATE, OMP_TARGET_ENTER_DATA, OMP_TARGET_EXIT_DATA stanza. * semantics.cc (cxx_omp_map_array_section): Avoid calling build_array_ref for non-array/non-pointer bases (error reported already). gcc/testsuite/ * c-c++-common/gomp/declare-mapper-15.c: New test. * c-c++-common/gomp/declare-mapper-16.c: New test. * g++.dg/gomp/declare-mapper-1.C: Adjust expected scan output. (cherry picked from commit 3f8c7483112807f8f2c48852f9cbad9a4d5b1da4) Diff: --- gcc/c-family/c-common.h | 2 +- gcc/c-family/c-omp.cc | 193 ++++++++++++++++++++- gcc/c/c-parser.cc | 11 +- gcc/c/c-tree.h | 1 - gcc/cp/parser.cc | 15 +- gcc/cp/pt.cc | 8 +- gcc/cp/semantics.cc | 5 +- .../c-c++-common/gomp/declare-mapper-15.c | 59 +++++++ .../c-c++-common/gomp/declare-mapper-16.c | 39 +++++ gcc/testsuite/g++.dg/gomp/declare-mapper-1.C | 2 +- 10 files changed, 313 insertions(+), 22 deletions(-) diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h index 89517e2a80cc..a4cd67861092 100644 --- a/gcc/c-family/c-common.h +++ b/gcc/c-family/c-common.h @@ -1357,7 +1357,7 @@ extern void c_omp_mark_declare_variant (location_t, tree, tree); extern void c_omp_adjust_map_clauses (tree, bool); template<typename T> struct omp_mapper_list; extern void c_omp_find_nested_mappers (struct omp_mapper_list<tree> *, tree); -extern tree c_omp_instantiate_mappers (tree); +extern tree c_omp_instantiate_mappers (tree, enum c_omp_region_type); namespace omp_addr_tokenizer { struct omp_addr_token; } typedef omp_addr_tokenizer::omp_addr_token omp_addr_token; diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc index bb121ede06ba..f75f0788ac66 100644 --- a/gcc/c-family/c-omp.cc +++ b/gcc/c-family/c-omp.cc @@ -4384,13 +4384,189 @@ remap_mapper_decl_1 (tree *tp, int *walk_subtrees, void *data) return NULL_TREE; } +static enum gomp_map_kind +omp_split_map_kind (enum gomp_map_kind op, bool *force_p, bool *always_p, + bool *present_p) +{ + *force_p = *always_p = *present_p = false; + + switch (op) + { + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_FORCE_PRESENT: + *force_p = true; + break; + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_ALWAYS_TOFROM: + *always_p = true; + break; + case GOMP_MAP_ALWAYS_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + *always_p = true; + /* Fallthrough. */ + case GOMP_MAP_PRESENT_ALLOC: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_PRESENT_TOFROM: + *present_p = true; + break; + default: + ; + } + + switch (op) + { + case GOMP_MAP_ALLOC: + case GOMP_MAP_FORCE_ALLOC: + case GOMP_MAP_PRESENT_ALLOC: + return GOMP_MAP_ALLOC; + case GOMP_MAP_TO: + case GOMP_MAP_FORCE_TO: + case GOMP_MAP_ALWAYS_TO: + case GOMP_MAP_PRESENT_TO: + case GOMP_MAP_ALWAYS_PRESENT_TO: + return GOMP_MAP_TO; + case GOMP_MAP_FROM: + case GOMP_MAP_FORCE_FROM: + case GOMP_MAP_ALWAYS_FROM: + case GOMP_MAP_PRESENT_FROM: + case GOMP_MAP_ALWAYS_PRESENT_FROM: + return GOMP_MAP_FROM; + case GOMP_MAP_TOFROM: + case GOMP_MAP_FORCE_TOFROM: + case GOMP_MAP_ALWAYS_TOFROM: + case GOMP_MAP_PRESENT_TOFROM: + case GOMP_MAP_ALWAYS_PRESENT_TOFROM: + return GOMP_MAP_TOFROM; + default: + ; + } + + return op; +} + +static enum gomp_map_kind +omp_join_map_kind (enum gomp_map_kind op, bool force_p, bool always_p, + bool present_p) +{ + gcc_assert (!force_p || !(always_p || present_p)); + + switch (op) + { + case GOMP_MAP_ALLOC: + if (force_p) + return GOMP_MAP_FORCE_ALLOC; + else if (present_p) + return GOMP_MAP_PRESENT_ALLOC; + break; + + case GOMP_MAP_TO: + if (force_p) + return GOMP_MAP_FORCE_TO; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TO; + else if (always_p) + return GOMP_MAP_ALWAYS_TO; + else if (present_p) + return GOMP_MAP_PRESENT_TO; + break; + + case GOMP_MAP_FROM: + if (force_p) + return GOMP_MAP_FORCE_FROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_FROM; + else if (always_p) + return GOMP_MAP_ALWAYS_FROM; + else if (present_p) + return GOMP_MAP_PRESENT_FROM; + break; + + case GOMP_MAP_TOFROM: + if (force_p) + return GOMP_MAP_FORCE_TOFROM; + else if (always_p && present_p) + return GOMP_MAP_ALWAYS_PRESENT_TOFROM; + else if (always_p) + return GOMP_MAP_ALWAYS_TOFROM; + else if (present_p) + return GOMP_MAP_PRESENT_TOFROM; + break; + + default: + ; + } + + return op; +} + +/* Map kind decay (OpenMP 5.2, 5.8.8 "declare mapper Directive"). Return the + map kind to use given MAPPER_KIND specified in the mapper and INVOKED_AS + specified on the clause that invokes the mapper. See also + fortran/trans-openmp.cc:omp_map_decayed_kind. */ + +static enum gomp_map_kind +omp_map_decayed_kind (enum gomp_map_kind mapper_kind, + enum gomp_map_kind invoked_as, bool exit_p) +{ + if (invoked_as == GOMP_MAP_RELEASE || invoked_as == GOMP_MAP_DELETE) + return invoked_as; + + bool force_p, always_p, present_p; + + invoked_as = omp_split_map_kind (invoked_as, &force_p, &always_p, &present_p); + gomp_map_kind decay_to; + + switch (mapper_kind) + { + case GOMP_MAP_ALLOC: + if (exit_p && invoked_as == GOMP_MAP_FROM) + decay_to = GOMP_MAP_RELEASE; + else + decay_to = GOMP_MAP_ALLOC; + break; + + case GOMP_MAP_TO: + if (invoked_as == GOMP_MAP_FROM) + decay_to = exit_p ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC; + else if (invoked_as == GOMP_MAP_ALLOC) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_TO; + break; + + case GOMP_MAP_FROM: + if (invoked_as == GOMP_MAP_ALLOC || invoked_as == GOMP_MAP_TO) + decay_to = GOMP_MAP_ALLOC; + else + decay_to = GOMP_MAP_FROM; + break; + + case GOMP_MAP_TOFROM: + case GOMP_MAP_UNSET: + decay_to = invoked_as; + break; + + default: + gcc_unreachable (); + } + + return omp_join_map_kind (decay_to, force_p, always_p, present_p); +} + /* Instantiate a mapper MAPPER for expression EXPR, adding new clauses to OUTLIST. OUTER_KIND is the mapping kind to use if not already specified in the mapper declaration. */ static tree * omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, - enum gomp_map_kind outer_kind) + enum gomp_map_kind outer_kind, + enum c_omp_region_type ort) { tree clauses = OMP_DECLARE_MAPPER_CLAUSES (mapper); tree dummy_var = OMP_DECLARE_MAPPER_DECL (mapper); @@ -4447,8 +4623,10 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, walk_tree (&unshared, remap_mapper_decl_1, &map_info, NULL); - if (OMP_CLAUSE_MAP_KIND (unshared) == GOMP_MAP_UNSET) - OMP_CLAUSE_SET_MAP_KIND (unshared, outer_kind); + enum gomp_map_kind decayed_kind + = omp_map_decayed_kind (clause_kind, outer_kind, + (ort & C_ORT_EXIT_DATA) != 0); + OMP_CLAUSE_SET_MAP_KIND (unshared, decayed_kind); type = TYPE_MAIN_VARIANT (type); @@ -4465,11 +4643,8 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); if (nested_mapper != mapper) { - if (clause_kind == GOMP_MAP_UNSET) - clause_kind = outer_kind; - outlist = omp_instantiate_mapper (outlist, nested_mapper, - t, clause_kind); + t, outer_kind, ort); continue; } } @@ -4491,7 +4666,7 @@ omp_instantiate_mapper (tree *outlist, tree mapper, tree expr, visible in the current parsing context. */ tree -c_omp_instantiate_mappers (tree clauses) +c_omp_instantiate_mappers (tree clauses, enum c_omp_region_type ort) { tree c, *pc, mapper_name = NULL_TREE; @@ -4564,7 +4739,7 @@ c_omp_instantiate_mappers (tree clauses) { tree mapper = lang_hooks.decls.omp_extract_mapper_directive (mapper_fn); - pc = omp_instantiate_mapper (pc, mapper, t, kind); + pc = omp_instantiate_mapper (pc, mapper, t, kind, ort); using_mapper = true; } else if (mapper_name) diff --git a/gcc/c/c-parser.cc b/gcc/c/c-parser.cc index 4cf2d8851662..9c0e65a74e5b 100644 --- a/gcc/c/c-parser.cc +++ b/gcc/c/c-parser.cc @@ -26921,7 +26921,9 @@ c_parser_omp_target_data (location_t loc, c_parser *parser, bool *if_p) tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, - "#pragma omp target data"); + "#pragma omp target data", false); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = c_finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -27079,7 +27081,9 @@ c_parser_omp_target_enter_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, - "#pragma omp target enter data"); + "#pragma omp target enter data", false); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = c_finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -27190,6 +27194,7 @@ c_parser_omp_target_exit_data (location_t loc, c_parser *parser, tree clauses = c_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", false); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA); clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; @@ -27449,7 +27454,7 @@ c_parser_omp_target (c_parser *parser, enum pragma_context context, bool *if_p) OMP_CLAUSE_CHAIN (nc) = OMP_CLAUSE_CHAIN (c); OMP_CLAUSE_CHAIN (c) = nc; } - clauses = c_omp_instantiate_mappers (clauses); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET); clauses = c_finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); diff --git a/gcc/c/c-tree.h b/gcc/c/c-tree.h index e151095c3a19..0b863cb7d1d0 100644 --- a/gcc/c/c-tree.h +++ b/gcc/c/c-tree.h @@ -1009,7 +1009,6 @@ extern tree c_check_omp_declare_reduction_r (tree *, int *, void *); extern tree c_omp_mapper_id (tree); extern tree c_omp_mapper_decl (tree); extern void c_omp_scan_mapper_bindings (location_t, tree *, tree); -extern tree c_omp_instantiate_mappers (tree); extern bool c_check_in_current_scope (tree); extern void c_pushtag (location_t, tree, tree); extern void c_bind (location_t, tree, bool); diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index d52f5c97039c..edbd3431643c 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -51556,7 +51556,10 @@ cp_parser_omp_target_data (cp_parser *parser, cp_token *pragma_tok, bool *if_p) tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_DATA_CLAUSE_MASK, - "#pragma omp target data", pragma_tok); + "#pragma omp target data", pragma_tok, false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -51671,7 +51674,11 @@ cp_parser_omp_target_enter_data (cp_parser *parser, cp_token *pragma_tok, tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_ENTER_DATA_CLAUSE_MASK, - "#pragma omp target enter data", pragma_tok); + "#pragma omp target enter data", pragma_tok, + false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP); + clauses = finish_omp_clauses (clauses, C_ORT_OMP); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; for (tree *pc = &clauses; *pc;) @@ -51788,6 +51795,8 @@ cp_parser_omp_target_exit_data (cp_parser *parser, cp_token *pragma_tok, = cp_parser_omp_all_clauses (parser, OMP_TARGET_EXIT_DATA_CLAUSE_MASK, "#pragma omp target exit data", pragma_tok, false); + if (!processing_template_decl) + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_EXIT_DATA); clauses = finish_omp_clauses (clauses, C_ORT_OMP_EXIT_DATA); c_omp_adjust_map_clauses (clauses, false); int map_seen = 0; @@ -52084,7 +52093,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok, OMP_CLAUSE_CHAIN (c) = nc; } if (!processing_template_decl) - clauses = c_omp_instantiate_mappers (clauses); + clauses = c_omp_instantiate_mappers (clauses, C_ORT_OMP_TARGET); clauses = finish_omp_clauses (clauses, C_ORT_OMP_TARGET); c_omp_adjust_map_clauses (clauses, true); diff --git a/gcc/cp/pt.cc b/gcc/cp/pt.cc index 3deff77c50b2..25982f02820c 100644 --- a/gcc/cp/pt.cc +++ b/gcc/cp/pt.cc @@ -18942,8 +18942,8 @@ tsubst_omp_clauses (tree clauses, enum c_omp_region_type ort, new_clauses = nreverse (new_clauses); if (ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_OMP_DECLARE_MAPPER) { - if (ort == C_ORT_OMP_TARGET) - new_clauses = c_omp_instantiate_mappers (new_clauses); + if (ort & C_ORT_OMP) + new_clauses = c_omp_instantiate_mappers (new_clauses, ort); new_clauses = finish_omp_clauses (new_clauses, ort); if (linear_no_step) for (nc = new_clauses; nc; nc = OMP_CLAUSE_CHAIN (nc)) @@ -20540,7 +20540,9 @@ tsubst_stmt (tree t, tree args, tsubst_flags_t complain, tree in_decl) case OMP_TARGET_UPDATE: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: - tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), C_ORT_OMP, args, + tmp = tsubst_omp_clauses (OMP_STANDALONE_CLAUSES (t), + (TREE_CODE (t) == OMP_TARGET_EXIT_DATA + ? C_ORT_OMP_EXIT_DATA : C_ORT_OMP), args, complain, in_decl); t = copy_node (t); OMP_STANDALONE_CLAUSES (t) = tmp; diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc index 18808354b6e8..6535376a5f97 100644 --- a/gcc/cp/semantics.cc +++ b/gcc/cp/semantics.cc @@ -6950,7 +6950,10 @@ cxx_omp_map_array_section (location_t loc, tree t) if (TREE_CODE (TREE_TYPE (t)) == REFERENCE_TYPE) t = convert_from_reference (t); - t = build_array_ref (loc, t, low); + if (TYPE_PTR_P (TREE_TYPE (t))) + t = build_array_ref (loc, t, low); + else + t = error_mark_node; } return t; diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c new file mode 100644 index 000000000000..ecda2e5ebd1a --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-15.c @@ -0,0 +1,59 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +typedef struct { + int a, b, c, d; +} S; + +int main () +{ + S s; + #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \ + map(from: x.c) map(tofrom: x.d) + + #pragma omp target enter data map(to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: 4\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(alloc: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(release: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(release:s\.c \[len: [0-9]+\]\) map\(release:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(present, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(force_present:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(present, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(force_present:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(always, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(always,to:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(always,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(always, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,from:s\.c \[len: [0-9]+\]\) map\(always,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + + #pragma omp target enter data map(always, present, to: s) + + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(force_present:s\.a \[len: [0-9]+\]\) map\(always,present,to:s\.b \[len: [0-9]+\]\) map\(force_present:s\.c \[len: [0-9]+\]\) map\(always,present,to:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + #pragma omp target exit data map(always, present, from: s) + + /* { dg-final { scan-tree-dump-times {map\(release:s\.a \[len: [0-9]+\]\) map\(release:s\.b \[len: [0-9]+\]\) map\(always,present,from:s\.c \[len: [0-9]+\]\) map\(always,present,from:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c new file mode 100644 index 000000000000..20383cc2d69f --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/declare-mapper-16.c @@ -0,0 +1,39 @@ +/* { dg-do compile } */ +/* { dg-options "-fopenmp -fdump-tree-gimple" } */ + +typedef struct { + int a, b, c, d; +} S; + +int main () +{ + S s = { 0, 0, 0, 0 }; + #pragma omp declare mapper (S x) map(alloc: x.a) map(to: x.b) \ + map(from: x.c) map(tofrom: x.d) + + #pragma omp target data map(s) + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(to:s\.b \[len: [0-9]+\]\) map\(from:s\.c \[len: [0-9]+\]\) map\(tofrom:s\.d \[len: [0-9]+\]\)} 3 "gimple" } } */ + { + #pragma omp target + { + s.a++; + s.b++; + s.c++; + s.d++; + } + } + + #pragma omp target data map(alloc: s) + /* { dg-final { scan-tree-dump-times {map\(struct:s \[len: 4\]\) map\(alloc:s\.a \[len: [0-9]+\]\) map\(alloc:s\.b \[len: [0-9]+\]\) map\(alloc:s\.c \[len: [0-9]+\]\) map\(alloc:s\.d \[len: [0-9]+\]\)} 1 "gimple" } } */ + { + #pragma omp target + { + s.a++; + s.b++; + s.c++; + s.d++; + } + } + + return 0; +} diff --git a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C index 2523d02b4a6f..2f2dd219bcb5 100644 --- a/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C +++ b/gcc/testsuite/g++.dg/gomp/declare-mapper-1.C @@ -55,4 +55,4 @@ int main (int argc, char *argv[]) } // { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(tofrom:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 4 "gimple" } } -// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } } +// { dg-final { scan-tree-dump-times {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(alloc:s\.size \[len: [0-9]+\]\) map\(alloc:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} 1 "gimple" } }
