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
