Re: [PATCH v2] OpenMP 5.0: Allow multiple clauses mapping same variable
Paul-Antoine Arras <[email protected]>
| Newsgroups | gmane.comp.gcc.patches,gmane.comp.gcc.fortran |
|---|---|
| Message-ID | <[email protected]> |
Forgot to mention in the commit message, but it fixes those OpenMP_VV tests: FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.F90 compile -O0 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.F90 compile -O3 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.F90 execute -O0 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.F90 execute -O3 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.c compile -O0 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.c compile -O3 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.c execute -O0 FAIL -> PASS: openmp_vv.sum:5.0 - test_target_mapping_before_alloc.c execute -O3 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_iterators.c compile -O0 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_iterators.c compile -O3 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_iterators.c execute -O0 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_iterators.c execute -O3 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_present_ordering.c compile -O0 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_present_ordering.c compile -O3 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_present_ordering.c execute -O0 FAIL -> PASS: openmp_vv.sum:5.1 - test_target_map_present_ordering.c execute -O3 On 15/05/2026 17:53, Paul-Antoine Arras wrote: > This patch allows multiple clauses on the same construct to map the same > variable, which was not valid in OpenMP 4.5, but allowed in 5.0. > > Internally, map clauses have to be deduplicated or merged before reaching the > topological sort in gimplify.cc, lest they might result in a cycle. This happens > in two places: first in the respective front-ends before any clause expansion, > then in the gimplifier just before grouping. The second pass is necessary due to > early clause expansion in the FE reintroducing some duplication (see > map-multi-2.f90). > > To make duplicate detection and folding easier in Fortran, enum gfc_omp_map_op > is adjusted to have the two least signficant bits mapped to FROM and TO, similar > to gomp_map_kind in gomp-constants.h > > gcc/c/ChangeLog: > > * c-typeck.cc (c_finish_omp_clauses): Call omp_remove_duplicate_maps > before clause expansion. > > gcc/cp/ChangeLog: > > * semantics.cc (finish_omp_clauses): Likewise. > > gcc/ChangeLog: > > * fold-const.cc (operand_compare::operand_equal_p): Handle > OMP_ARRAY_SECTION. > * gimplify.cc (gimplify_scan_omp_clauses): Call > omp_remove_duplicate_maps after partial clause expansion. > * omp-general.cc (omp_remove_duplicate_maps): New function. > * omp-general.h (omp_remove_duplicate_maps): Declare. > * omp-low.cc (install_var_field): Add new 'tree key_expr = NULL_TREE' > default parameter. Set splay-tree lookup key to key_expr instead of > var if key_expr is non-NULL. Adjust call to install_parm_decl. > Update comments. > (scan_sharing_clauses): Use clause tree expression as splay-tree key > for map/to/from and OpenACC firstprivate cases when installing the > variable field into the send/receive record type. > (lower_oacc_reductions): Adjust to find map-clause of reduction > variable, then create receiver-ref. > (lower_omp_target): Adjust to lookup var field using clause expression. > > gcc/fortran/ChangeLog: > > * gfortran.h (enum gfc_omp_map_op): Dedicate the two LSB to TO and FROM. > * openmp.cc (resolve_omp_clauses): Adjust to allow duplicate > mapped variables for OpenMP. > * trans-openmp.cc (gfc_trans_omp_clauses): Remove duplicates before > clause expansion. > > libgomp/ChangeLog: > > * testsuite/libgomp.c++/target-map-multi-1.C: New test. > * testsuite/libgomp.c-c++-common/target-map-iterators-6.c: New test. > * testsuite/libgomp.c-c++-common/target-map-multi-1.c: New test. > * testsuite/libgomp.c-c++-common/target-map-multi-2.c: New test. > * testsuite/libgomp.c-c++-common/target-map-multi-3.c: New test. > * testsuite/libgomp.c-c++-common/target-map-multi-4.c: New test. > * testsuite/libgomp.fortran/target-map-multi-1.f90: New test. > * testsuite/libgomp.fortran/target-map-multi-2.f90: New test. > * testsuite/libgomp.fortran/target-map-multi-3.f90: New test. > * testsuite/libgomp.fortran/target-map-multi-4.f90: New test. > * testsuite/libgomp.fortran/target-map-multi-5.f90: New test. > > gcc/testsuite/ChangeLog: > > * c-c++-common/gomp/clauses-2.c: Adjust testcase. > * c-c++-common/gomp/map-6.c: Adjust testcase. > * gfortran.dg/gomp/pr107214.f90: Adjust testcase. > * c-c++-common/gomp/map-multi-1.c: New test. > * c-c++-common/gomp/map-multi-2.c: New test. > * gfortran.dg/gomp/map-multi-1.f90: New test. > * gfortran.dg/gomp/map-multi-2.f90: New test. > > Co-Authored-By: Chung-Lin Tang <[email protected]> > Co-Authored-By: Sandra Loosemore <[email protected]> > --- > gcc/c/c-typeck.cc | 4 +- > gcc/cp/semantics.cc | 7 +- > gcc/fold-const.cc | 3 + > gcc/fortran/gfortran.h | 56 +++---- > gcc/fortran/openmp.cc | 2 +- > gcc/fortran/trans-openmp.cc | 25 +++ > gcc/gimplify.cc | 1 + > gcc/omp-general.cc | 90 +++++++++++ > gcc/omp-general.h | 2 + > gcc/omp-low.cc | 82 +++++++--- > gcc/testsuite/c-c++-common/gomp/clauses-2.c | 4 +- > gcc/testsuite/c-c++-common/gomp/map-6.c | 4 +- > gcc/testsuite/c-c++-common/gomp/map-multi-1.c | 36 +++++ > gcc/testsuite/c-c++-common/gomp/map-multi-2.c | 41 +++++ > .../gfortran.dg/gomp/map-multi-1.f90 | 50 ++++++ > .../gfortran.dg/gomp/map-multi-2.f90 | 60 ++++++++ > gcc/testsuite/gfortran.dg/gomp/pr107214.f90 | 5 +- > .../libgomp.c++/target-map-multi-1.C | 113 ++++++++++++++ > .../target-map-iterators-6.c | 43 ++++++ > .../libgomp.c-c++-common/target-map-multi-1.c | 72 +++++++++ > .../libgomp.c-c++-common/target-map-multi-2.c | 54 +++++++ > .../libgomp.c-c++-common/target-map-multi-3.c | 22 +++ > .../libgomp.c-c++-common/target-map-multi-4.c | 66 ++++++++ > .../libgomp.fortran/target-map-multi-1.f90 | 83 ++++++++++ > .../libgomp.fortran/target-map-multi-2.f90 | 84 +++++++++++ > .../libgomp.fortran/target-map-multi-3.f90 | 18 +++ > .../libgomp.fortran/target-map-multi-4.f90 | 87 +++++++++++ > .../libgomp.fortran/target-map-multi-5.f90 | 142 ++++++++++++++++++ > 28 files changed, 1193 insertions(+), 63 deletions(-) > create mode 100644 gcc/testsuite/c-c++-common/gomp/map-multi-1.c > create mode 100644 gcc/testsuite/c-c++-common/gomp/map-multi-2.c > create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90 > create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90 > create mode 100644 libgomp/testsuite/libgomp.c++/target-map-multi-1.C > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c > create mode 100644 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c > create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90 > create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90 > create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90 > create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90 > create mode 100644 libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90 > > diff --git a/gcc/c/c-typeck.cc b/gcc/c/c-typeck.cc > index 043e49794ca..08553724e54 100644 > --- a/gcc/c/c-typeck.cc > +++ b/gcc/c/c-typeck.cc > @@ -16643,6 +16643,8 @@ c_oacc_check_attachments (tree c) > tree > c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > { > + clauses = omp_remove_duplicate_maps (clauses, true); > + > bitmap_head generic_head, firstprivate_head, lastprivate_head; > bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; > bitmap_head oacc_reduction_head, is_on_device_head; > @@ -17872,7 +17874,7 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > } > else if (bitmap_bit_p (&map_head, DECL_UID (t)) > && !bitmap_bit_p (&map_field_head, DECL_UID (t)) > - && ort != C_ORT_OMP > + && ort != C_ORT_OMP && ort != C_ORT_OMP_TARGET > && ort != C_ORT_OMP_EXIT_DATA) > { > if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) > diff --git a/gcc/cp/semantics.cc b/gcc/cp/semantics.cc > index cd42695aa82..f1fbec0be49 100644 > --- a/gcc/cp/semantics.cc > +++ b/gcc/cp/semantics.cc > @@ -7796,6 +7796,8 @@ cp_finish_omp_init_prefer_type (tree pref_type) > tree > finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > { > + clauses = omp_remove_duplicate_maps (clauses, true); > + > bitmap_head generic_head, firstprivate_head, lastprivate_head; > bitmap_head aligned_head, map_head, map_field_head, map_firstprivate_head; > bitmap_head oacc_reduction_head, is_on_device_head; > @@ -9564,8 +9566,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > == GOMP_MAP_FIRSTPRIVATE_POINTER)) > { > if (bitmap_bit_p (&generic_head, DECL_UID (t)) > - || bitmap_bit_p (&firstprivate_head, DECL_UID (t)) > - || bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))) > + || bitmap_bit_p (&firstprivate_head, DECL_UID (t))) > { > error_at (OMP_CLAUSE_LOCATION (c), > "%qD appears more than once in data clauses", t); > @@ -9588,7 +9589,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort) > bitmap_set_bit (&map_firstprivate_head, DECL_UID (t)); > else if (bitmap_bit_p (&map_head, DECL_UID (t)) > && !bitmap_bit_p (&map_field_head, DECL_UID (t)) > - && ort != C_ORT_OMP > + && ort != C_ORT_OMP && ort != C_ORT_OMP_TARGET > && ort != C_ORT_OMP_EXIT_DATA) > { > if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP) > diff --git a/gcc/fold-const.cc b/gcc/fold-const.cc > index 56503e570bc..0e5e3dd30de 100644 > --- a/gcc/fold-const.cc > +++ b/gcc/fold-const.cc > @@ -3816,6 +3816,9 @@ operand_compare::operand_equal_p (tree type0, const_tree arg0, > } > return false; > > + case OMP_ARRAY_SECTION: > + return OP_SAME (0) && OP_SAME (1) && OP_SAME (2); > + > default: > return false; > } > diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h > index b0ce54e1c21..aabae6cd554 100644 > --- a/gcc/fortran/gfortran.h > +++ b/gcc/fortran/gfortran.h > @@ -1347,33 +1347,33 @@ enum gfc_omp_depend_doacross_op > > enum gfc_omp_map_op > { > - OMP_MAP_ALLOC, > - OMP_MAP_IF_PRESENT, > - OMP_MAP_ATTACH, > - OMP_MAP_TO, > - OMP_MAP_FROM, > - OMP_MAP_TOFROM, > - OMP_MAP_DELETE, > - OMP_MAP_DETACH, > - OMP_MAP_FORCE_ALLOC, > - OMP_MAP_FORCE_TO, > - OMP_MAP_FORCE_FROM, > - OMP_MAP_FORCE_TOFROM, > - OMP_MAP_FORCE_PRESENT, > - OMP_MAP_FORCE_DEVICEPTR, > - OMP_MAP_DEVICE_RESIDENT, > - OMP_MAP_LINK, > - OMP_MAP_RELEASE, > - OMP_MAP_ALWAYS_TO, > - OMP_MAP_ALWAYS_FROM, > - OMP_MAP_ALWAYS_TOFROM, > - OMP_MAP_PRESENT_ALLOC, > - OMP_MAP_PRESENT_TO, > - OMP_MAP_PRESENT_FROM, > - OMP_MAP_PRESENT_TOFROM, > - OMP_MAP_ALWAYS_PRESENT_TO, > - OMP_MAP_ALWAYS_PRESENT_FROM, > - OMP_MAP_ALWAYS_PRESENT_TOFROM > + OMP_MAP_ALLOC = 0, > + OMP_MAP_TO = 1 << 0, > + OMP_MAP_FROM = 1 << 1, > + OMP_MAP_TOFROM = OMP_MAP_TO | OMP_MAP_FROM, > + OMP_MAP_IF_PRESENT = 1 << 2, > + OMP_MAP_ATTACH = 1 << 3, > + OMP_MAP_DELETE = 1 << 4, > + OMP_MAP_DETACH = 1 << 5, > + OMP_MAP_FORCE_ALLOC = 1 << 6, > + OMP_MAP_FORCE_TO = OMP_MAP_FORCE_ALLOC | OMP_MAP_TO, > + OMP_MAP_FORCE_FROM = OMP_MAP_FORCE_ALLOC | OMP_MAP_FROM, > + OMP_MAP_FORCE_TOFROM = OMP_MAP_FORCE_ALLOC | OMP_MAP_TOFROM, > + OMP_MAP_FORCE_PRESENT = 1 << 7, > + OMP_MAP_FORCE_DEVICEPTR = 1 << 8, > + OMP_MAP_DEVICE_RESIDENT = 1 << 9, > + OMP_MAP_LINK = 1 << 10, > + OMP_MAP_RELEASE = 1 << 11, > + OMP_MAP_ALWAYS_TO = (1 << 12) | OMP_MAP_TO, > + OMP_MAP_ALWAYS_FROM = (1 << 12) | OMP_MAP_FROM, > + OMP_MAP_ALWAYS_TOFROM = (1 << 12) | OMP_MAP_TOFROM, > + OMP_MAP_PRESENT_ALLOC = 1 << 13, > + OMP_MAP_PRESENT_TO = (1 << 13) | OMP_MAP_TO, > + OMP_MAP_PRESENT_FROM = (1 << 13) | OMP_MAP_FROM, > + OMP_MAP_PRESENT_TOFROM = (1 << 13) | OMP_MAP_TOFROM, > + OMP_MAP_ALWAYS_PRESENT_TO = OMP_MAP_ALWAYS_TO | OMP_MAP_PRESENT_TO, > + OMP_MAP_ALWAYS_PRESENT_FROM = OMP_MAP_ALWAYS_FROM | OMP_MAP_PRESENT_FROM, > + OMP_MAP_ALWAYS_PRESENT_TOFROM = OMP_MAP_ALWAYS_TOFROM | OMP_MAP_PRESENT_TOFROM > }; > > enum gfc_omp_defaultmap > @@ -1421,7 +1421,7 @@ typedef struct gfc_omp_namelist > gfc_omp_depend_doacross_op depend_doacross_op; > struct > { > - ENUM_BITFIELD (gfc_omp_map_op) op:8; > + ENUM_BITFIELD (gfc_omp_map_op) op : 16; > bool readonly; > } map; > gfc_expr *align; > diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc > index ee3bcb3ba63..8b6545533c6 100644 > --- a/gcc/fortran/openmp.cc > +++ b/gcc/fortran/openmp.cc > @@ -9296,7 +9296,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, > gfc_error ("Symbol %qs has mixed component and non-component " > "accesses at %L", n->sym->name, &n->where); > } > - else if (n->sym->mark) > + else if (list != OMP_LIST_MAP && n->sym->mark) > gfc_error ("Symbol %qs present on multiple clauses at %L", > n->sym->name, &n->where); > else > diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc > index 538ba045530..393c6527e1e 100644 > --- a/gcc/fortran/trans-openmp.cc > +++ b/gcc/fortran/trans-openmp.cc > @@ -4283,6 +4283,31 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, > case OMP_LIST_MAP: > for (; n != NULL; n = n->next) > { > + // Remove duplicates > + bool skip = false; > + for (gfc_omp_namelist *n2 = n->next; n2 != NULL; n2 = n2->next) > + { > + if (n2->sym == n->sym > + && gfc_dep_compare_expr (n2->expr, n->expr) == 0) > + { > + if (n2->u.map.op == n->u.map.op) > + { > + skip = true; > + break; > + } > + else if ((n2->u.map.op & ~OMP_MAP_TOFROM) > + == (n->u.map.op & ~OMP_MAP_TOFROM)) > + { > + n2->u.map.op = (enum gfc_omp_map_op) (n->u.map.op > + | n2->u.map.op); > + skip = true; > + break; > + } > + } > + } > + if (skip) > + continue; > + > if (!n->sym->attr.referenced > || n->sym->attr.flavor == FL_PARAMETER) > continue; > diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc > index 8bfc71315f9..76bc4458251 100644 > --- a/gcc/gimplify.cc > +++ b/gcc/gimplify.cc > @@ -13817,6 +13817,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, > || code == OACC_UPDATE > || code == OACC_DECLARE) > { > + *list_p = omp_remove_duplicate_maps (*list_p, false); > groups = omp_gather_mapping_groups (list_p); > > if (groups) > diff --git a/gcc/omp-general.cc b/gcc/omp-general.cc > index 11f6917f7dd..ca293c920e8 100644 > --- a/gcc/omp-general.cc > +++ b/gcc/omp-general.cc > @@ -5151,3 +5151,93 @@ omp_merge_context_selectors (location_t loc, tree outer_ctx, tree inner_ctx, > merged_ctx = nreverse (merged_ctx); > return omp_check_context_selector (loc, merged_ctx, directive); > } > + > +/* Remove duplicate and merge clauses mapping the same variable. This function > + is called twice: FIRST in the C and C++ front-ends before any clause > + expansion happens, then in the gimplifier before gathering groups. This is > + because it is easier to process most clauses earlier but some duplicates > + still get introduced during the early clause expansion in the front-ends. */ > + > +tree > +omp_remove_duplicate_maps (tree clauses, bool first) > +{ > + if (clauses == NULL_TREE) > + return NULL_TREE; > + > + tree outlist = NULL_TREE; > + tree *outlist_p = &outlist; > + bool remove = false; > + tree c1; > + for (c1 = clauses; OMP_CLAUSE_CHAIN (c1) != NULL_TREE; > + c1 = OMP_CLAUSE_CHAIN (c1)) > + { > + if (OMP_CLAUSE_CODE (c1) != OMP_CLAUSE_MAP) > + { > + *outlist_p = c1; > + outlist_p = &OMP_CLAUSE_CHAIN (*outlist_p); > + continue; > + } > + > + for (tree c2 = OMP_CLAUSE_CHAIN (c1); c2 != NULL_TREE; > + c2 = OMP_CLAUSE_CHAIN (c2)) > + { > + if (OMP_CLAUSE_CODE (c2) != OMP_CLAUSE_MAP) > + continue; > + > + bool maybe_dup_found > + = (OMP_CLAUSE_CODE (c1) == OMP_CLAUSE_CODE (c2) > + && ((OMP_CLAUSE_DECL (c1) == NULL_TREE > + && OMP_CLAUSE_DECL (c2) == NULL_TREE) > + || operand_equal_p (OMP_CLAUSE_DECL (c1), > + OMP_CLAUSE_DECL (c2))) > + && ((OMP_CLAUSE_SIZE (c1) == NULL_TREE > + && OMP_CLAUSE_SIZE (c2) == NULL_TREE) > + || (OMP_CLAUSE_SIZE (c1) != NULL_TREE > + && OMP_CLAUSE_SIZE (c2) != NULL_TREE > + && operand_equal_p (OMP_CLAUSE_SIZE (c1), > + OMP_CLAUSE_SIZE (c2)))) > + && ((OMP_CLAUSE_ITERATORS (c1) == NULL_TREE > + && OMP_CLAUSE_ITERATORS (c2) == NULL_TREE) > + || operand_equal_p (OMP_CLAUSE_ITERATORS (c1), > + OMP_CLAUSE_ITERATORS (c2)))); > + if (maybe_dup_found) > + { > + if (first) > + { > + if (OMP_CLAUSE_MAP_KIND (c1) == OMP_CLAUSE_MAP_KIND (c2)) > + { > + remove = true; > + break; > + } > + else if ((OMP_CLAUSE_MAP_KIND (c1) & ~GOMP_MAP_TOFROM) > + == (OMP_CLAUSE_MAP_KIND (c2) & ~GOMP_MAP_TOFROM)) > + { > + OMP_CLAUSE_SET_MAP_KIND (c2, > + (OMP_CLAUSE_MAP_KIND (c1) > + | OMP_CLAUSE_MAP_KIND (c2))); > + remove = true; > + break; > + } > + } > + /* When called from the gimplifier, only compare bits up to > + GOMP_MAP_FLAG_SPECIAL_2. Higher flags may need to be present > + multiple times.*/ > + else if (OMP_CLAUSE_MAP_KIND (c1) == OMP_CLAUSE_MAP_KIND (c2) > + && (OMP_CLAUSE_MAP_KIND (c1) & ~0b11111) == 0) > + { > + remove = true; > + break; > + } > + } > + } > + if (remove) > + remove = false; > + else > + { > + *outlist_p = c1; > + outlist_p = &OMP_CLAUSE_CHAIN (*outlist_p); > + } > + } > + *outlist_p = c1; > + return outlist; > +} > diff --git a/gcc/omp-general.h b/gcc/omp-general.h > index 4a6d58471de..a5126269650 100644 > --- a/gcc/omp-general.h > +++ b/gcc/omp-general.h > @@ -417,4 +417,6 @@ extern bool omp_parse_expr (vec<omp_addr_token *> &, tree); > extern tree omp_loop_number_of_iterations (tree, int, tree * = NULL); > extern void omp_maybe_apply_loop_xforms (tree *, tree); > > +extern tree omp_remove_duplicate_maps (tree, bool); > + > #endif /* GCC_OMP_GENERAL_H */ > diff --git a/gcc/omp-low.cc b/gcc/omp-low.cc > index b93012107f1..466ffb35454 100644 > --- a/gcc/omp-low.cc > +++ b/gcc/omp-low.cc > @@ -784,24 +784,40 @@ build_sender_ref (tree var, omp_context *ctx) > return build_sender_ref ((splay_tree_key) var, ctx); > } > > -/* Add a new field for VAR inside the structure CTX->SENDER_DECL. If > - BASE_POINTERS_RESTRICT, declare the field with restrict. */ > +/* Add a new field for VAR inside the structure CTX. If BY_REF is true, > + use a pointer to the VAR rather than VAR itself. > + MASK is a bit mask of other options. Bits are interpreted as: > + 1: Install VAR in ctx->field_map. > + 2: Install VAR in ctx->sfield_map. > + 4: VAR is an array, convert it to a pointer. > + 8: Use DECL_UID (VAR) instead of VAR as key. > + 16: Use DECL_NAME (VAR) instead of VAR as key. > + 32: Don't dereference omp_is_reference types. > + KEY_EXPR allows specifying something other than VAR as the lookup key. > + If specified, it overrides the 8 and 16 MASK bits. */ > > static void > -install_var_field (tree var, bool by_ref, int mask, omp_context *ctx) > +install_var_field (tree var, bool by_ref, int mask, omp_context *ctx, > + tree key_expr = NULL_TREE) > { > tree field, type, sfield = NULL_TREE; > splay_tree_key key = (splay_tree_key) var; > > - if ((mask & 16) != 0) > + if (key_expr) > + /* Allow caller to explicitly set the expression used as the key. */ > + key = (splay_tree_key) key_expr; > + else > { > - key = (splay_tree_key) &DECL_NAME (var); > - gcc_checking_assert (key != (splay_tree_key) var); > - } > - if ((mask & 8) != 0) > - { > - key = (splay_tree_key) &DECL_UID (var); > - gcc_checking_assert (key != (splay_tree_key) var); > + if ((mask & 16) != 0) > + { > + key = (splay_tree_key) &DECL_NAME (var); > + gcc_checking_assert (key != (splay_tree_key) var); > + } > + if ((mask & 8) != 0) > + { > + key = (splay_tree_key) &DECL_UID (var); > + gcc_checking_assert (key != (splay_tree_key) var); > + } > } > gcc_assert ((mask & 1) == 0 > || !splay_tree_lookup (ctx->field_map, key)); > @@ -1389,8 +1405,13 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > || (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR > && lang_hooks.decls.omp_array_data (decl, true))) > { > + /* OpenACC firstprivate clauses are later processed with same > + code path as map clauses in lower_omp_target, so follow > + the same convention of using the whole clause expression > + as splay-tree key. */ > + tree k = (is_oacc_parallel_or_serial (ctx) ? c : NULL_TREE); > by_ref = !omp_privatize_by_reference (decl); > - install_var_field (decl, by_ref, 3, ctx); > + install_var_field (decl, by_ref, 3, ctx, k); > } > else if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_HAS_DEVICE_ADDR) > { > @@ -1683,7 +1704,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > gcc_assert (INDIRECT_REF_P (decl2)); > decl2 = TREE_OPERAND (decl2, 0); > gcc_assert (DECL_P (decl2)); > - install_var_field (decl2, true, 3, ctx); > + install_var_field (decl2, true, 3, ctx, c); > install_var_local (decl2, ctx); > install_var_local (decl, ctx); > } > @@ -1693,9 +1714,9 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER > && !OMP_CLAUSE_MAP_ZERO_BIAS_ARRAY_SECTION (c) > && TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE) > - install_var_field (decl, true, 7, ctx); > + install_var_field (decl, true, 7, ctx, c); > else > - install_var_field (decl, true, 3, ctx); > + install_var_field (decl, true, 3, ctx, c); > if (is_gimple_omp_offloaded (ctx->stmt) > && !(is_gimple_omp_oacc (ctx->stmt) > && OMP_CLAUSE_MAP_IN_REDUCTION (c))) > @@ -1730,7 +1751,7 @@ scan_sharing_clauses (tree clauses, omp_context *ctx) > FIELD_DECL, NULL_TREE, ptr_type_node); > SET_DECL_ALIGN (field, TYPE_ALIGN (ptr_type_node)); > insert_field_into_struct (ctx->record_type, field); > - splay_tree_insert (ctx->field_map, (splay_tree_key) decl, > + splay_tree_insert (ctx->field_map, (splay_tree_key) c, > (splay_tree_value) field); > } > } > @@ -7453,6 +7474,7 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, > gcc_checking_assert (!is_oacc_kernels_decomposed_part (ctx)); > > tree orig = OMP_CLAUSE_DECL (c); > + tree orig_clause; > tree var = maybe_lookup_decl (orig, ctx); > tree ref_to_res = NULL_TREE; > tree incoming, outgoing, v1, v2, v3; > @@ -7523,10 +7545,20 @@ lower_oacc_reductions (location_t loc, tree clauses, tree level, bool inner, > do_lookup: > /* This is the outermost construct with this reduction, > see if there's a mapping for it. */ > - if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET > - && maybe_lookup_field (orig, outer) && !is_private) > + orig_clause = NULL_TREE; > + if (gimple_code (outer->stmt) == GIMPLE_OMP_TARGET) > + for (tree cls = gimple_omp_target_clauses (outer->stmt); > + cls; cls = OMP_CLAUSE_CHAIN (cls)) > + if (OMP_CLAUSE_CODE (cls) == OMP_CLAUSE_MAP > + && orig == OMP_CLAUSE_DECL (cls) > + && maybe_lookup_field (cls, outer)) > + { > + orig_clause = cls; > + break; > + } > + if (orig_clause != NULL_TREE && !is_private) > { > - ref_to_res = build_receiver_ref (orig, false, outer); > + ref_to_res = build_receiver_ref (orig_clause, false, outer); > if (omp_privatize_by_reference (orig)) > ref_to_res = build_simple_mem_ref (ref_to_res); > > @@ -12949,7 +12981,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > continue; > } > > - if (!maybe_lookup_field (var, ctx)) > + if (!maybe_lookup_field (c, ctx)) > continue; > > /* Don't remap compute constructs' reduction variables, because the > @@ -12958,7 +12990,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > && is_gimple_omp_oacc (ctx->stmt) > && OMP_CLAUSE_MAP_IN_REDUCTION (c))) > { > - x = build_receiver_ref (var, true, ctx); > + x = build_receiver_ref (c, true, ctx); > tree new_var = lookup_decl (var, ctx); > > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > @@ -13353,7 +13385,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > } > else > { > - tree x = build_sender_ref (ovar, ctx); > + tree x = build_sender_ref (c, ctx); > tree v = ovar; > if (in_reduction_clauses > && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > @@ -13402,7 +13434,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > gcc_assert (DECL_P (ovar2)); > ovar = ovar2; > } > - if (!maybe_lookup_field (ovar, ctx) > + if (!maybe_lookup_field (c, ctx) > && !(OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH > || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))) > @@ -13452,7 +13484,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > } > else if (nc) > { > - x = build_sender_ref (ovar, ctx); > + x = build_sender_ref (nc, ctx); > > if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP > && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER > @@ -14416,7 +14448,7 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context *ctx) > type = TREE_TYPE (type); > ref_to_ptr = true; > } > - x = build_receiver_ref (OMP_CLAUSE_DECL (prev), false, ctx); > + x = build_receiver_ref (prev, false, ctx); > x = fold_convert_loc (clause_loc, type, x); > if (!integer_zerop (OMP_CLAUSE_SIZE (c))) > { > diff --git a/gcc/testsuite/c-c++-common/gomp/clauses-2.c b/gcc/testsuite/c-c++-common/gomp/clauses-2.c > index 8f98d57a312..d5018fa1ffa 100644 > --- a/gcc/testsuite/c-c++-common/gomp/clauses-2.c > +++ b/gcc/testsuite/c-c++-common/gomp/clauses-2.c > @@ -15,9 +15,9 @@ foo (int *p, int q, struct S t, int i, int j, int k, int l) > bar (p); > #pragma omp target map (p) , map (p[0]) > bar (p); > - #pragma omp target map (q) map (q) /* { dg-error "appears more than once in map clauses" } */ > + #pragma omp target map(q) map(q) > bar (&q); > - #pragma omp target map (p[0]) map (p[0]) /* { dg-error "appears more than once in data clauses" } */ > + #pragma omp target map(p[0]) map(p[0]) > bar (p); > #pragma omp target map (t) map (t.r) > bar (&t.r); > diff --git a/gcc/testsuite/c-c++-common/gomp/map-6.c b/gcc/testsuite/c-c++-common/gomp/map-6.c > index c1f3d1079b4..aa3e9786e5c 100644 > --- a/gcc/testsuite/c-c++-common/gomp/map-6.c > +++ b/gcc/testsuite/c-c++-common/gomp/map-6.c > @@ -157,10 +157,10 @@ foo (void) > #pragma omp target map (always, close) > ; > > - #pragma omp target map (always, always) /* { dg-error "'always' appears more than once in map clauses" } */ > + #pragma omp target map (always, always) > ; > > - #pragma omp target map (always, always, close) /* { dg-error "'always' appears more than once in map clauses" } */ > + #pragma omp target map (always, always, close) > ; > > #pragma omp target map (always, close, to: always, close, b7) > diff --git a/gcc/testsuite/c-c++-common/gomp/map-multi-1.c b/gcc/testsuite/c-c++-common/gomp/map-multi-1.c > new file mode 100644 > index 00000000000..496dcf6c77c > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/gomp/map-multi-1.c > @@ -0,0 +1,36 @@ > +/* { dg-do compile } */ > +/* { dg-additional-options "-fdump-tree-original" } */ > + > +/* Check that extraneous clauses mapping the same variable multiple times are > + either removed or merged. */ > + > +#define DIM 17 > + > +void f (int *x) > +{ > + #pragma omp target map(alloc: x) map(to: x) map(alloc: x) map(from: x) map(alloc: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(alloc: x) map(alloc: x) map(alloc: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(alloc: x) map(to: x) map(alloc: x) map(to: x) map(alloc: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(alloc: x) map(from: x) map(alloc: x) map(from: x) map(alloc: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > +} > + > +/* { dg-final { scan-tree-dump-times {map\(tofrom:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(alloc:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(to:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(from:x\)} 1 "original" } } */ > diff --git a/gcc/testsuite/c-c++-common/gomp/map-multi-2.c b/gcc/testsuite/c-c++-common/gomp/map-multi-2.c > new file mode 100644 > index 00000000000..3eb4c1c004a > --- /dev/null > +++ b/gcc/testsuite/c-c++-common/gomp/map-multi-2.c > @@ -0,0 +1,41 @@ > +/* { dg-do compile } */ > +/* { dg-additional-options "-fdump-tree-original" } */ > + > +/* Check that map clauses are only merged if they have the same modifiers. */ > + > +#define DIM 17 > + > +void f (int *x) > +{ > + #pragma omp target map(always, to: x) map(tofrom: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(tofrom: x) map(always, from: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(to: x) map(present, alloc: x) map(from: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(close,always,from: x) map(close,present,alloc: x) map(present,to: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > + #pragma omp target map(always,to: x) map(always,present,alloc: x) map(present,from: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i] += i; > + } > +} > + > +/* { dg-final { scan-tree-dump-times {map\(tofrom:x\) map\(always,to:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(always,from:x\) map\(tofrom:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(present,alloc:x\) map\(tofrom:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(present,to:x\) map\(always,from:x\)} 1 "original" } } */ > +/* { dg-final { scan-tree-dump-times {map\(present,from:x\) map\(always,to:x\)} 1 "original" } } */ > diff --git a/gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90 b/gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90 > new file mode 100644 > index 00000000000..c12b6d431cf > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90 > @@ -0,0 +1,50 @@ > +! { dg-do compile } > +! { dg-additional-options "-fdump-tree-original" } > + > +! Check that extraneous clauses mapping the same variable multiple times are > +! either removed or merged. > + > +subroutine f > + implicit none > + integer :: x, i > + > + !$omp target map(alloc: x) map(to: x) map(alloc: x) map(from: x) map(alloc: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(alloc: x) map(alloc: x) map(alloc: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(alloc: x) map(to: x) map(alloc: x) map(to: x) map(alloc: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(alloc: x) map(from: x) map(alloc: x) map(from: x) map(alloc: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > +end subroutine f > + > +! { dg-final { scan-tree-dump-times {map\(tofrom:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(alloc:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(to:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(from:x\)} 1 "original" } } > + > +subroutine g (a) > + integer, intent(inout) :: a(:) > + > + !$omp target map(to: a) map(alloc: a) map(from: a) > + a = a * 4 > + !$omp end target > + !$omp target map(alloc: a) map(tofrom: a) map(alloc: a) > + a = a * 8 > + !$omp end target > + > +end subroutine > + > +! { dg-final { scan-tree-dump-times {map\(tofrom:\*a\.0\) map\(alloc:a\.0 \[pointer assign, bias: 0\]\)} 2 "original" } } > diff --git a/gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90 b/gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90 > new file mode 100644 > index 00000000000..fbd4dcbb256 > --- /dev/null > +++ b/gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90 > @@ -0,0 +1,60 @@ > +! { dg-do compile } > +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } > + > +! Check that map clauses are only merged if they have the same modifiers. > +! Also check that the GIMPLE pass removes duplicates resulting from early clause > +! expansion in the front-end. > + > +subroutine f > + implicit none > + integer :: x, i > + > + !$omp target map(always,to: x) map(tofrom: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(tofrom: x) map(always,from: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(to: x) map(present,alloc: x) map(from: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(close,always,from: x) map(close,present,alloc: x) map(present,to: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > + !$omp target map(always,to: x) map(always,present,alloc: x) map(present,from: x) > + do i = 1, 17 > + x = x + i > + end do > + !$omp end target > +end subroutine f > + > +! { dg-final { scan-tree-dump-times {map\(always,to:x\) map\(tofrom:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(tofrom:x\) map\(always,from:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(present,alloc:x\) map\(tofrom:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(always,from:x\) map\(present,to:x\)} 1 "original" } } > +! { dg-final { scan-tree-dump-times {map\(always,to:x\) map\(present,from:x\)} 1 "original" } } > + > +subroutine g (a) > + integer, intent(inout) :: a(:) > + > + !$omp target map(always, to: a) map(tofrom: a) > + a = a * 4 > + !$omp end target > + !$omp target map(close, present, tofrom: a) map(always, from: a) > + a = a * 8 > + !$omp end target > + > +end subroutine > + > +! { dg-final { scan-tree-dump-times {map\(alloc:a\.0 \[pointer assign, bias: 0\]\)} 4 "original" } } > +! { dg-final { scan-tree-dump-times {map\(alloc:a\.0 \[pointer assign, bias: 0\]\)} 2 "gimple" } } > + > + > diff --git a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 > index 25949934e84..84141b43727 100644 > --- a/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 > +++ b/gcc/testsuite/gfortran.dg/gomp/pr107214.f90 > @@ -1,7 +1,10 @@ > ! { dg-do compile } > +! { dg-additional-options "-fdump-tree-original" } > > program p > integer, allocatable :: a > - !$omp target map(tofrom: a, a) ! { dg-error "Symbol 'a' present on multiple clauses" } > + !$omp target map(tofrom: a, a) > !$omp end target > end > + > +! { dg-final { scan-tree-dump-times {map\(tofrom:\*a\)} 1 "original" } } > diff --git a/libgomp/testsuite/libgomp.c++/target-map-multi-1.C b/libgomp/testsuite/libgomp.c++/target-map-multi-1.C > new file mode 100644 > index 00000000000..54507176cf0 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c++/target-map-multi-1.C > @@ -0,0 +1,113 @@ > +// Test multiple map clauses for the same variable using pointer array sections > +// and C++ references. > +// { dg-do run } > + > +extern "C" void abort (); > + > +static const int N = 100; > + > +/* Test fixed-size array with multiple map clauses. */ > +static void > +test_pointer () > +{ > + int p[N]; > + > + /* map(to) + map(alloc) + map(from): three clauses on the same array. */ > + for (int i = 0; i < N; i++) p[i] = i; > + #pragma omp target map(to: p) map(alloc: p) map(from: p) > + for (int i = 0; i < N; i++) p[i] *= 4; > + for (int i = 0; i < N; i++) if (p[i] != i * 4) abort (); > + > + /* map(to) + map(tofrom): tofrom covers both directions. */ > + for (int i = 0; i < N; i++) p[i] = i; > + #pragma omp target map(to: p) map(tofrom: p) > + for (int i = 0; i < N; i++) p[i] *= 5; > + for (int i = 0; i < N; i++) if (p[i] != i * 5) abort (); > + > + /* map(alloc) + map(to): device gets host values via 'to'. */ > + for (int i = 0; i < N; i++) p[i] = i; > + int sum = 0; > + #pragma omp target map(alloc: p) map(to: p) map(tofrom: sum) > + for (int i = 0; i < N; i++) sum += p[i]; > + if (sum != N * (N - 1) / 2) abort (); > + > + /* map(alloc) + map(from): device values come back via 'from'. */ > + #pragma omp target map(alloc: p) map(from: p) > + for (int i = 0; i < N; i++) p[i] = i * 7; > + for (int i = 0; i < N; i++) if (p[i] != i * 7) abort (); > + > + /* map(alloc) + map(tofrom) + map(alloc): three clauses, full bidirectional. */ > + for (int i = 0; i < N; i++) p[i] = i; > + #pragma omp target map(alloc: p) map(tofrom: p) map(alloc: p) > + for (int i = 0; i < N; i++) p[i] *= 8; > + for (int i = 0; i < N; i++) if (p[i] != i * 8) abort (); > +} > + > +/* Test C++ array references: r aliases a. */ > +static void > +test_reference () > +{ > + int a[N]; > + int (&r)[N] = a; > + > + /* map(to: r) + map(from: r): send and receive via the same reference. */ > + for (int i = 0; i < N; i++) a[i] = i; > + #pragma omp target map(to: r) map(from: r) > + for (int i = 0; i < N; i++) r[i] *= 4; > + for (int i = 0; i < N; i++) if (r[i] != i * 4) abort (); > + > + /* map(to: r) + map(tofrom: r): to and tofrom on the same reference. */ > + for (int i = 0; i < N; i++) a[i] = i; > + #pragma omp target map(to: r) map(tofrom: r) > + for (int i = 0; i < N; i++) r[i] *= 5; > + for (int i = 0; i < N; i++) if (r[i] != i * 5) abort (); > + > + /* map(to: r) + map(alloc: r) + map(from: r): three clauses on the same reference. */ > + for (int i = 0; i < N; i++) a[i] = i; > + #pragma omp target map(to: r) map(alloc: r) map(from: r) > + for (int i = 0; i < N; i++) r[i] *= 6; > + for (int i = 0; i < N; i++) if (r[i] != i * 6) abort (); > + > + /* map(alloc: r) + map(to: r): alloc and to on the same reference. */ > + for (int i = 0; i < N; i++) a[i] = i; > + int sum = 0; > + #pragma omp target map(alloc: r) map(to: r) map(tofrom: sum) > + for (int i = 0; i < N; i++) sum += r[i]; > + if (sum != N * (N - 1) / 2) abort (); > +} > + > +/* Test pointer + array sections with enter/exit data using multiple map > + clauses. */ > +static void > +test_pointer_enter_exit () > +{ > + int *p = new int[N]; > + > + /* map(alloc) + map(to) on enter data. */ > + for (int i = 0; i < N; i++) p[i] = i; > + #pragma omp target enter data map(alloc: p[0:N]) map(to: p[0:N]) > + int sum = 0; > + #pragma omp target map(alloc: p[0:N]) map(tofrom: sum) > + for (int i = 0; i < N; i++) sum += p[i]; > + if (sum != N * (N - 1) / 2) abort (); > + #pragma omp target exit data map(delete: p[0:N]) > + > + /* map(release) + map(from) on exit data: copy back then release. */ > + for (int i = 0; i < N; i++) p[i] = i; > + #pragma omp target enter data map(to: p[0:N]) > + #pragma omp target map(alloc: p[0:N]) > + for (int i = 0; i < N; i++) p[i] *= 3; > + #pragma omp target exit data map(release: p[0:N]) map(from: p[0:N]) > + for (int i = 0; i < N; i++) if (p[i] != i * 3) abort (); > + > + delete[] p; > +} > + > +int > +main () > +{ > + test_pointer (); > + test_reference (); > + test_pointer_enter_exit (); > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c > new file mode 100644 > index 00000000000..b6013a99ac9 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c > @@ -0,0 +1,43 @@ > +/* { dg-do run } */ > +/* { dg-require-effective-target offload_device_nonshared_as } */ > + > +/* This testcase used to ICE. Test that an array can appear more than once > + in map clauses in combination with iterators. */ > + > +#include <stdlib.h> > + > +#define DIM 128 > + > +void > +make_array (int *x[], int dim) > +{ > + for (int i = 0; i < DIM; i++) > + { > + x[i] = (int *) malloc (sizeof (int)); > + *(x[i]) = i; > + } > +} > + > +int > +check_array (int *x[], int dim) > +{ > + for (int i = 0; i < DIM; i++) > + if (*(x[i]) != -i) > + return 1; > + return 0; > +} > + > +int > +main (void) > +{ > + int *x[DIM]; > + make_array (x, DIM); > + > +#pragma omp target map (iterator(it = 0:DIM), tofrom: x[it][:1]) map (to: x) > + { > + for (int i = 0; i < DIM; i++) > + x[i][0] = -x[i][0]; > + } > + > + return check_array (x, DIM); > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c > new file mode 100644 > index 00000000000..6d4d4c31f72 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c > @@ -0,0 +1,72 @@ > +/* Test multiple map clauses for the same variable with various > + combinations of map-types: alloc, to, from, tofrom. */ > +/* { dg-do run } */ > + > +#define N 100 > + > +int > +main (void) > +{ > + int a[N]; > + int sum; > + > + /* map(to) + map(from) = map(tofrom). */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target map(to: a) map(alloc: a) map(from: a) > + for (int i = 0; i < N; i++) > + a[i] *= 4; > + for (int i = 0; i < N; i++) > + if (a[i] != i * 4) > + __builtin_abort (); > + > + /* map(to) + map(tofrom): tofrom covers both directions. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target map(to: a) map(tofrom: a) map(to: a) > + for (int i = 0; i < N; i++) > + a[i] *= 5; > + for (int i = 0; i < N; i++) > + if (a[i] != i * 5) > + __builtin_abort (); > + > + /* map(from) + map(tofrom): tofrom covers both directions. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target map(from: a) map(tofrom: a) map(from: a) > + for (int i = 0; i < N; i++) > + a[i] *= 6; > + for (int i = 0; i < N; i++) > + if (a[i] != i * 6) > + __builtin_abort (); > + > + /* map(alloc) + map(to): device gets host values via 'to'. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + sum = 0; > + #pragma omp target map(alloc: a) map(to: a) map(tofrom: sum) map(alloc: sum) > + for (int i = 0; i < N; i++) > + sum += a[i]; > + if (sum != N * (N - 1) / 2) > + __builtin_abort (); > + > + /* map(alloc) + map(from): device values come back to host via 'from'. */ > + #pragma omp target map(alloc: a) map(from: a) map(alloc: a) > + for (int i = 0; i < N; i++) > + a[i] = i * 7; > + for (int i = 0; i < N; i++) > + if (a[i] != i * 7) > + __builtin_abort (); > + > + /* map(alloc) + map(tofrom): full bidirectional transfer. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target map(alloc: a) map(tofrom: a) map(alloc: a) > + for (int i = 0; i < N; i++) > + a[i] *= 8; > + for (int i = 0; i < N; i++) > + if (a[i] != i * 8) > + __builtin_abort (); > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c > new file mode 100644 > index 00000000000..bb72dfcab49 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c > @@ -0,0 +1,54 @@ > +/* Test multiple map clauses for the same variable with the 'always' > + map-type modifier. */ > +/* { dg-do run } */ > + > +#define N 100 > + > +int > +main (void) > +{ > + int a[N]; > + > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target data map(tofrom: a) > + { > + /* Device has a[i] = i. Update the host copy to create divergence. */ > + for (int i = 0; i < N; i++) > + a[i] = i * 2; > + > + /* map(always, to) + map(tofrom): 'always' forces the updated host > + values (i*2) onto the device despite the existing mapping. */ > + #pragma omp target map(always, to: a) map(tofrom: a) > + { > + for (int i = 0; i < N; i++) > + if (a[i] != i * 2) > + __builtin_abort (); > + for (int i = 0; i < N; i++) > + a[i] = i * 3; > + } > + > + for (int i = 0; i < N; i++) > + if (a[i] != i * 2) > + __builtin_abort (); > + > + /* Reset host for the next test; device retains i*3. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + > + /* map(tofrom) + map(always, from): 'always' forces new device values > + back to host. */ > + #pragma omp target map(tofrom: a) map(always, from: a) > + for (int i = 0; i < N; i++) > + { > + if (a[i] != i * 3) > + __builtin_abort (); > + a[i] = i * 4; > + } > + for (int i = 0; i < N; i++) > + if (a[i] != i * 4) > + __builtin_abort (); > + } > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c > new file mode 100644 > index 00000000000..c7aac627382 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c > @@ -0,0 +1,22 @@ > +/* Test multiple map clauses for the same variable with the 'present' > + map-type modifier. */ > +/* { dg-do run } */ > + > +#define N 100 > + > +int > +main (void) > +{ > + int a[N]; > + > + for (int i = 0; i < N; i++) > + a[i] = i; > + > + /* { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } */ > + /* { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } */ > + #pragma omp target map(to: a) map(present, alloc: a) map(from: a) > + for (int i = 0; i < N; i++) > + a[i] *= 2; > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c > new file mode 100644 > index 00000000000..080ca93a88d > --- /dev/null > +++ b/libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c > @@ -0,0 +1,66 @@ > +/* Test multiple map clauses for the same variable in target enter/exit data > + constructs, including release and delete map-types. */ > +/* { dg-do run } */ > + > +#define N 100 > + > +int > +main (void) > +{ > + int a[N]; > + int sum; > + > + /* delete + release */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target enter data map(alloc: a) map(to: a) > + sum = 0; > + #pragma omp target map(alloc: a) map(tofrom: sum) > + for (int i = 0; i < N; i++) > + sum += a[i]; > + if (sum != N * (N - 1) / 2) > + __builtin_abort (); > + #pragma omp target exit data map(delete: a) map(release: a) > + > + > + /* release + release: duplicate release > + decrements the reference count once (deduplicated). */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target enter data map(to: a) /* refcount = 1 */ > + #pragma omp target enter data map(to: a) /* refcount = 2 */ > + #pragma omp target exit data map(release: a) map(release: a) /* refcount = 1 */ > + sum = 0; > + #pragma omp target map(alloc: a) map(tofrom: sum) > + for (int i = 0; i < N; i++) > + sum += a[i]; > + if (sum != N * (N - 1) / 2) > + __builtin_abort (); > + #pragma omp target exit data map(delete: a) /* refcount = 0 */ > + > + /* delete + delete: duplicate delete > + removes the mapping unconditionally once (deduplicated). */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target enter data map(to: a) > + #pragma omp target exit data map(delete: a) map(delete: a) > + > + /* from + release: copy device values back > + to host and release the mapping. */ > + for (int i = 0; i < N; i++) > + a[i] = i; > + #pragma omp target enter data map(to: a) /* refcount = 1 */ > + #pragma omp target > + for (int i = 0; i < N; i++) > + a[i] *= 3; > + #pragma omp target exit data map(release: a) map(from: a) /* refcount = 0 */ > + for (int i = 0; i < N; i++) > + if (a[i] != i * 3) > + __builtin_abort (); > + > + for (int i = 0; i < N; i++) > + if (a[i] != i * 3) > + __builtin_abort (); > + > + return 0; > +} > diff --git a/libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90 b/libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90 > new file mode 100644 > index 00000000000..0dc41bb136a > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90 > @@ -0,0 +1,83 @@ > +! Test multiple map clauses for the same variable with various > +! combinations of map-types: alloc, to, from, tofrom. > +! { dg-do run } > + > +program main > + implicit none > + integer, parameter :: N = 100 > + integer :: a(N), s, i > + > + ! map(to) + map(alloc) + map(from) = map(tofrom). > + a = [(i, i=1,N)] > + !$omp target map(to: a) map(alloc: a) map(from: a) > + a = a * 4 > + !$omp end target > + if (any (a /= [(i*4, i=1,N)])) stop 1 > + > + ! map(to) + map(tofrom) + map(to): tofrom covers both directions. > + a = [(i, i=1,N)] > + !$omp target map(to: a) map(tofrom: a) map(to: a) > + a = a * 5 > + !$omp end target > + if (any (a /= [(i*5, i=1,N)])) stop 2 > + > + ! map(from) + map(tofrom) + map(from): tofrom covers both directions. > + a = [(i, i=1,N)] > + !$omp target map(from: a) map(tofrom: a) map(from: a) > + a = a * 6 > + !$omp end target > + if (any (a /= [(i*6, i=1,N)])) stop 3 > + > + ! map(alloc) + map(to): device gets host values via 'to'. > + ! Also tests map(tofrom) + map(alloc) on the scalar s. > + a = [(i, i=1,N)] > + s = 0 > + !$omp target map(alloc: a) map(to: a) map(tofrom: s) map(alloc: s) > + do i = 1, N > + s = s + a(i) > + end do > + !$omp end target > + if (s /= N * (N + 1) / 2) stop 4 > + > + ! map(alloc) + map(from) + map(alloc): device values come back via 'from'. > + !$omp target map(alloc: a) map(from: a) map(alloc: a) > + do i = 1, N > + a(i) = i * 7 > + end do > + !$omp end target > + if (any (a /= [(i*7, i=1,N)])) stop 5 > + > + ! map(alloc) + map(tofrom) + map(alloc): full bidirectional transfer. > + a = [(i, i=1,N)] > + !$omp target map(alloc: a) map(tofrom: a) map(alloc: a) > + a = a * 8 > + !$omp end target > + if (any (a /= [(i*8, i=1,N)])) stop 6 > + > + ! Same tests via a subroutine to verify correct behaviour with > + ! dummy arguments (passed by reference). > + call test_dummy (a) > + > +contains > + > + subroutine test_dummy (a) > + integer, intent(inout) :: a(:) > + integer :: i > + > + ! map(to) + map(alloc) + map(from) with dummy argument. > + a = [(i, i=1,size(a))] > + !$omp target map(to: a) map(alloc: a) map(from: a) > + a = a * 4 > + !$omp end target > + if (any (a /= [(i*4, i=1,size(a))])) stop 7 > + > + ! map(alloc) + map(tofrom) + map(alloc) with dummy argument. > + a = [(i, i=1,size(a))] > + !$omp target map(alloc: a) map(tofrom: a) map(alloc: a) > + a = a * 8 > + !$omp end target > + if (any (a /= [(i*8, i=1,size(a))])) stop 8 > + > + end subroutine > + > +end program > diff --git a/libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90 b/libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90 > new file mode 100644 > index 00000000000..783a7aa1f58 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90 > @@ -0,0 +1,84 @@ > +! Test multiple map clauses for the same variable with the 'always' > +! map-type modifier. > +! { dg-do run } > + > +program main > + implicit none > + integer, parameter :: N = 100 > + integer :: a(N), i > + > + a = [(i, i=1,N)] > + !$omp target data map(tofrom: a) > + ! Device has a(i) = i. Update the host copy to create divergence. > + a = [(i*2, i=1,N)] > + > + ! map(always, to) + map(tofrom): 'always' forces the updated host > + ! values (i*2) onto the device despite the existing mapping. > + !$omp target map(always, to: a) map(tofrom: a) > + do i = 1, N > + if (a(i) /= i*2) stop 1 > + end do > + a = [(i*3, i=1,N)] > + !$omp end target > + > + do i = 1, N > + if (a(i) /= i*2) stop 2 > + end do > + > + ! Reset host for the next test; device retains i*3. > + a = [(i, i=1,N)] > + > + ! map(tofrom) + map(always, from): 'always' forces new device values > + ! back to host. > + !$omp target map(tofrom: a) map(always, from: a) > + do i = 1, N > + if (a(i) /= i*3) stop 3 > + a(i) = i*4 > + end do > + !$omp end target > + do i = 1, N > + if (a(i) /= i*4) stop 4 > + end do > + !$omp end target data > + > + ! Same tests via a subroutine to verify correct behaviour with > + ! dummy arguments (passed by reference). > + call test_dummy (a) > + > +contains > + > + subroutine test_dummy (a) > + integer, intent(inout) :: a(:) > + integer :: i > + > + a = [(i, i=1,size(a))] > + !$omp target data map(tofrom: a) > + a = [(i*2, i=1,size(a))] > + > + !$omp target map(always, to: a) map(tofrom: a) > + do i = 1, size(a) > + if (a(i) /= i*2) stop 5 > + end do > + a = [(i*3, i=1,size(a))] > + !$omp end target > + > + do i = 1, size(a) > + if (a(i) /= i*2) stop 6 > + end do > + > + a = [(i, i=1,size(a))] > + > + !$omp target map(tofrom: a) map(always, from: a) > + do i = 1, size(a) > + if (a(i) /= i*3) stop 7 > + a(i) = i*4 > + end do > + !$omp end target > + do i = 1, size(a) > + if (a(i) /= i*4) stop 8 > + end do > + !$omp end target data > + > + end subroutine > + > +end program > diff --git a/libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90 b/libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90 > new file mode 100644 > index 00000000000..44b601853d9 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90 > @@ -0,0 +1,18 @@ > +! Test multiple map clauses for the same variable with the 'present' > +! modifier. > +! { dg-do run } > + > +program main > + implicit none > + integer, parameter :: N = 100 > + integer :: a(N), i > + > + a = [(i, i=1,N)] > + > + ! { dg-output "libgomp: present clause: not present on the device \\(addr: 0x\[0-9a-f\]+, size: \[0-9\]+ \\(0x\[0-9a-f\]+\\), dev: \[0-9\]+\\\)" { target offload_device_nonshared_as } } > + ! { dg-shouldfail "present error triggered" { offload_device_nonshared_as } } > + !$omp target map(to: a) map(present, alloc: a) map(from: a) > + a = a * 2 > + !$omp end target > + > +end program > diff --git a/libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90 b/libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90 > new file mode 100644 > index 00000000000..88a14913611 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90 > @@ -0,0 +1,87 @@ > +! Test multiple map clauses for the same variable in target enter/exit > +! data constructs, including release and delete map-types. > +! { dg-do run } > + > +program main > + implicit none > + integer, parameter :: N = 100 > + integer :: a(N), s, i > + > + ! delete + release in exit data. > + a = [(i, i=1,N)] > + !$omp target enter data map(alloc: a) map(to: a) > + s = 0 > + !$omp target map(alloc: a) map(tofrom: s) > + do i = 1, N > + s = s + a(i) > + end do > + !$omp end target > + if (s /= N * (N + 1) / 2) stop 1 > + !$omp target exit data map(delete: a) map(release: a) > + > + ! release + release: duplicate release decrements the reference count > + ! once (deduplicated). Two enter data calls set refcount to 2, so after > + ! one deduplicated release refcount is 1 and the mapping remains. > + a = [(i, i=1,N)] > + !$omp target enter data map(to: a) ! refcount = 1 > + !$omp target enter data map(to: a) ! refcount = 2 > + !$omp target exit data map(release: a) map(release: a) ! refcount = 1 > + s = 0 > + !$omp target map(alloc: a) map(tofrom: s) > + do i = 1, N > + s = s + a(i) > + end do > + !$omp end target > + if (s /= N * (N + 1) / 2) stop 2 > + !$omp target exit data map(delete: a) ! refcount = 0 > + > + ! delete + delete: duplicate delete removes the mapping unconditionally > + ! once (deduplicated). > + a = [(i, i=1,N)] > + !$omp target enter data map(to: a) > + !$omp target exit data map(delete: a) map(delete: a) > + > + ! from + release: copy device values back to host and release the mapping. > + a = [(i, i=1,N)] > + !$omp target enter data map(to: a) ! refcount = 1 > + !$omp target > + a = a * 3 > + !$omp end target > + !$omp target exit data map(release: a) map(from: a) ! refcount = 0 > + if (any (a /= [(i*3, i=1,N)])) stop 3 > + if (any (a /= [(i*3, i=1,N)])) stop 4 > + > + ! Same tests via a subroutine to verify correct behaviour with > + ! dummy arguments (passed by reference). > + call test_dummy (a) > + > +contains > + > + subroutine test_dummy (a) > + integer, intent(inout) :: a(:) > + integer :: s, i > + > + ! delete + release with dummy argument. > + a = [(i, i=1,size(a))] > + !$omp target enter data map(alloc: a) map(to: a) > + s = 0 > + !$omp target map(alloc: a) map(tofrom: s) > + do i = 1, size(a) > + s = s + a(i) > + end do > + !$omp end target > + if (s /= size(a) * (size(a) + 1) / 2) stop 5 > + !$omp target exit data map(delete: a) map(release: a) > + > + ! from + release with dummy argument. > + a = [(i, i=1,size(a))] > + !$omp target enter data map(to: a) > + !$omp target > + a = a * 3 > + !$omp end target > + !$omp target exit data map(release: a) map(from: a) > + if (any (a /= [(i*3, i=1,size(a))])) stop 6 > + > + end subroutine > + > +end program > diff --git a/libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90 b/libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90 > new file mode 100644 > index 00000000000..94c7c941ed2 > --- /dev/null > +++ b/libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90 > @@ -0,0 +1,142 @@ > +! Test multiple map clauses for the same variable using Fortran array section > +! notation (subscript triplets): fixed-size, allocatable, and dummy arrays. > +! { dg-do run } > + > +program main > + implicit none > + integer, parameter :: N = 100 > + integer :: a(N), s, i > + integer, allocatable :: b(:) > + > + ! --- Fixed-size array --- > + > + ! map(to) + map(alloc) + map(from) on the same section. > + a = [(i, i=1,N)] > + !$omp target map(to: a(1:N)) map(alloc: a(1:N)) map(from: a(1:N)) > + a = a * 4 > + !$omp end target > + if (any (a /= [(i*4, i=1,N)])) stop 1 > + > + ! map(to) + map(tofrom) on the same section. > + a = [(i, i=1,N)] > + !$omp target map(to: a(1:N)) map(tofrom: a(1:N)) > + a = a * 5 > + !$omp end target > + if (any (a /= [(i*5, i=1,N)])) stop 2 > + > + ! map(alloc) + map(to): device gets host values via 'to'. > + a = [(i, i=1,N)] > + s = 0 > + !$omp target map(alloc: a(1:N)) map(to: a(1:N)) map(tofrom: s) > + do i = 1, N > + s = s + a(i) > + end do > + !$omp end target > + if (s /= N*(N+1)/2) stop 3 > + > + ! map(alloc) + map(from): device values come back via 'from'. > + !$omp target map(alloc: a(1:N)) map(from: a(1:N)) > + do i = 1, N > + a(i) = i * 7 > + end do > + !$omp end target > + if (any (a /= [(i*7, i=1,N)])) stop 4 > + > + ! map(alloc) + map(tofrom) + map(alloc): full bidirectional. > + a = [(i, i=1,N)] > + !$omp target map(alloc: a(1:N)) map(tofrom: a(1:N)) map(alloc: a(1:N)) > + a = a * 8 > + !$omp end target > + if (any (a /= [(i*8, i=1,N)])) stop 5 > + > + ! --- Allocatable array --- > + > + allocate (b(N)) > + > + ! map(to) + map(alloc) + map(from) on the same section. > + b = [(i, i=1,N)] > + !$omp target map(to: b(1:N)) map(alloc: b(1:N)) map(from: b(1:N)) > + b = b * 4 > + !$omp end target > + if (any (b /= [(i*4, i=1,N)])) stop 6 > + > + ! map(to) + map(tofrom) on the same section. > + b = [(i, i=1,N)] > + !$omp target map(to: b(1:N)) map(tofrom: b(1:N)) > + b = b * 5 > + !$omp end target > + if (any (b /= [(i*5, i=1,N)])) stop 7 > + > + ! map(alloc) + map(to): device gets host values via 'to'. > + b = [(i, i=1,N)] > + s = 0 > + !$omp target map(alloc: b(1:N)) map(to: b(1:N)) map(tofrom: s) > + do i = 1, N > + s = s + b(i) > + end do > + !$omp end target > + if (s /= N*(N+1)/2) stop 8 > + > + ! map(alloc) + map(from): device values come back via 'from'. > + !$omp target map(alloc: b(1:N)) map(from: b(1:N)) > + do i = 1, N > + b(i) = i * 7 > + end do > + !$omp end target > + if (any (b /= [(i*7, i=1,N)])) stop 9 > + > + ! map(alloc) + map(tofrom) + map(alloc): full bidirectional. > + b = [(i, i=1,N)] > + !$omp target map(alloc: b(1:N)) map(tofrom: b(1:N)) map(alloc: b(1:N)) > + b = b * 8 > + !$omp end target > + if (any (b /= [(i*8, i=1,N)])) stop 10 > + > + deallocate (b) > + > + ! Same tests via a subroutine to verify correct behaviour with > + ! dummy arguments (passed by reference / array descriptor). > + call test_dummy (a) > + > +contains > + > + subroutine test_dummy (a) > + integer, intent(inout) :: a(:) > + integer :: i, s, n > + > + n = size(a) > + > + ! map(to) + map(alloc) + map(from) on the same section. > + a = [(i, i=1,n)] > + !$omp target map(to: a(1:n)) map(alloc: a(1:n)) map(from: a(1:n)) > + a = a * 4 > + !$omp end target > + if (any (a /= [(i*4, i=1,n)])) stop 11 > + > + ! map(to) + map(tofrom) on the same section. > + a = [(i, i=1,n)] > + !$omp target map(to: a(1:n)) map(tofrom: a(1:n)) > + a = a * 5 > + !$omp end target > + if (any (a /= [(i*5, i=1,n)])) stop 12 > + > + ! map(alloc) + map(to): device gets host values via 'to'. > + a = [(i, i=1,n)] > + s = 0 > + !$omp target map(alloc: a(1:n)) map(to: a(1:n)) map(tofrom: s) > + do i = 1, n > + s = s + a(i) > + end do > + !$omp end target > + if (s /= n*(n+1)/2) stop 13 > + > + ! map(alloc) + map(tofrom) + map(alloc): full bidirectional. > + a = [(i, i=1,n)] > + !$omp target map(alloc: a(1:n)) map(tofrom: a(1:n)) map(alloc: a(1:n)) > + a = a * 8 > + !$omp end target > + if (any (a /= [(i*8, i=1,n)])) stop 14 > + > + end subroutine > + > +end program -- PA