Thanks for the swift review!
See updated patch in attachment and comments inline below.
On 15/05/2026 21:19, Tobias Burnus wrote:
@Thomas – this code also affects OpenACC. Thus, you might want
to have a look as well.
I think we have two options: Make this OpenMP only or
also use this for OpenACC code. (See also second comment
below.) - WDYT?
* * *
...
* * *
[BUG]
I think we need to file a PR for the following issue (unchanged by your
patch);
that's C and C++ (Fortran is fine):
void f() {
int x, y;
#pragma omp target map(tofrom:x,y) private(y)
;
#pragma omp target map(tofrom:x,y) firstprivate(x)
;
}
gives:
test.c:5:31: error: ‘x’ appears both in data and map clauses
5 | #pragma omp target map(tofrom:x,y) firstprivate(x)
That's correct – but IMHO it should give the very same error also for
PRIVATE.
Filed PR: https://gcc.gnu.org/PR125367
* * *
[OpenACC]
void f() {
int X;
#pragma acc data copyin(X) copyout(X)
X += 5;
}
Before this patch:
foo.c:3:26: error: ‘X’ appears more than once in data clauses
3 | #pragma acc data copyin(X) copyout(X)
After the patch:
a-foo.c.006t.original: #pragma acc data map(tofrom:X)
a-foo.c.007t.gimple: #pragma omp target oacc_data map(tofrom:X
[len: 4])
a-foo.c.010t.omplower: #pragma omp target oacc_data map(tofrom:X
[len: 4])
I believe that's fine - also light of the still open OpenACC Issue #552
(+ pull req. draft #554). It seems as if OpenACC indented to permit
specifying the same variable multiple times (at least it does not
disallow it). With this patch, the following option is handled: Comment
from August last year (listening multiple options): "If we only want to
handle perfect aliasing, then sorting clauses at
compile time is probably sufficient"
If more should be supported, doing the simple cases at compile time is
fine - leaving the rest to the runtime. Otherwise, it is invalid, but
likely usually not detected because it is unlikely that one implements
an expensive check in the runtime.
Thus, we should be fine.
HOWEVER: I guess, we need some some testcases for this OpenACC case?
Alternatively, we have to/could exclude it from applying to OpenACC.
NOTE: Contrary to OpenMP, there is currently no testcase in this
regard at all - neither one checking that merging workes by this
patch nor any testcase that checks for 'Error: Symbol ‘x’ present
on multiple clauses' (which existed before for OpenMP and are
now removed.)
I must say I have overlooked the OpenACC case. In the current state of
the patch, a number of such testcases are actually failing (e.g.
testsuite/c-c++-common/goacc/data-clause-duplicate-1.c.). If it is
decided that this patch should indeed apply to OpenACC as well, I'll
have to update them.
* * *
--- a/gcc/c/c-typeck.cc
+++ b/gcc/c/c-typeck.cc
@@ -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)
{
I think the coding style at least in that FE is to have the '&&'
aligned, i.e. C_ORT_OMP_TARGET on a line of its own.
* * *
--- a/gcc/cp/semantics.cc
+++ b/gcc/cp/semantics.cc
@@ -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)
{
Fixed alignment in both FE.
--- a/gcc/omp-general.cc
+++ b/gcc/omp-general.cc
...
+tree
+omp_remove_duplicate_maps (tree clauses, bool first)
+ 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)))
I believe this cannot occur; we have sometimes NULL_POINTER for 'omp_…'
special identifiers and we might have some error conditions, but I doubt
we generate any map entry for those – and having an error_mark_node
seems to be more likely if something goes wrong later.
However regarding:
+ && ((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))))
The NULL_TREE seems to occur mostly when called from the front end as
the size is not yet filled in. (It might also be the case during the
gimplification, I have not checked what comes first.)
+ && ((OMP_CLAUSE_ITERATORS (c1) == NULL_TREE
+ && OMP_CLAUSE_ITERATORS (c2) == NULL_TREE)
+ || operand_equal_p (OMP_CLAUSE_ITERATORS (c1),
+ OMP_CLAUSE_ITERATORS (c2))));
while this one is obvious.
Maybe a comment makes sense that the OMP_CLAUSE_DECL is just to be
prudent but
it shouldn't occur while OMP_CLAUSE_SIZE often gets filled in only later?
Added comments for the cases where OMP_CLAUSE_DECL or OMP_CLAUSE_SIZE
are NULL_TREE.
(In any case, it is now documented here.)
* * *
+ if (first)
…
+ /* When called from the gimplifier, only compare bits up to
+ GOMP_MAP_FLAG_SPECIAL_2. Higher flags may need to be present
+ multiple times.*/
This might be clearer if worded as
When called from the gimplifier, remove duplicate map clauses with
identical kind only when the bits above GOMP_MAP_FLAG_SPECIAL_2 are
unset - as clauses with those flags set may need to be present
multiple times.
The current wording it not really wrong, but I first read it backward.
Rephrased comment as suggested.
--- 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. */
I would write 'If specified, it also overrides ...' - I think the 'also'
makes it clearer as by default the tree 'VAR' is used as key.
Added "also" to the comment.
* * *
LGTM – with the style/comments remarks addressed or at least considered.
Please give Thomas a few days (like early next week) to comment on OpenACC.
Thanks,
--
PA
commit 73082b44712d2f6ae5f9c715167ba548d672022a
Author: Paul-Antoine Arras <[email protected]>
Date: Fri May 15 17:45:01 2026 +0200
OpenMP 5.0: Allow multiple clauses mapping same variable
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]>
diff --git gcc/c/c-typeck.cc gcc/c/c-typeck.cc
index f36bf539b68..ed071ed6275 100644
--- gcc/c/c-typeck.cc
+++ gcc/c/c-typeck.cc
@@ -16684,6 +16684,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;
@@ -17914,6 +17916,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_TARGET
&& ort != C_ORT_OMP_EXIT_DATA)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
diff --git gcc/cp/semantics.cc gcc/cp/semantics.cc
index 6564d9e37a6..3fb63a8f7a0 100644
--- gcc/cp/semantics.cc
+++ gcc/cp/semantics.cc
@@ -7843,6 +7843,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;
@@ -9611,8 +9613,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);
@@ -9636,6 +9637,7 @@ 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_TARGET
&& ort != C_ORT_OMP_EXIT_DATA)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
diff --git gcc/fold-const.cc gcc/fold-const.cc
index a7351445eac..41e8043ac9b 100644
--- gcc/fold-const.cc
+++ 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 gcc/fortran/gfortran.h gcc/fortran/gfortran.h
index 37a8582e36d..16fc5e52cd9 100644
--- gcc/fortran/gfortran.h
+++ 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 gcc/fortran/openmp.cc gcc/fortran/openmp.cc
index a90600952f1..354d3ab4744 100644
--- gcc/fortran/openmp.cc
+++ gcc/fortran/openmp.cc
@@ -9293,7 +9293,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 gcc/fortran/trans-openmp.cc gcc/fortran/trans-openmp.cc
index 680e43c1cc0..b36b64be161 100644
--- gcc/fortran/trans-openmp.cc
+++ gcc/fortran/trans-openmp.cc
@@ -4308,6 +4308,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 gcc/gimplify.cc gcc/gimplify.cc
index e4db4b1d9bd..51e9605db87 100644
--- gcc/gimplify.cc
+++ gcc/gimplify.cc
@@ -13836,6 +13836,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 gcc/omp-general.cc gcc/omp-general.cc
index cf8cddc50bc..7ab9770b98f 100644
--- gcc/omp-general.cc
+++ gcc/omp-general.cc
@@ -5152,3 +5152,98 @@ 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)
+ && ((/* In the current state, a map clause decl is not supposed
+ to be NULL; but let's be defensive. */
+ OMP_CLAUSE_DECL (c1) == NULL_TREE
+ && OMP_CLAUSE_DECL (c2) == NULL_TREE)
+ || operand_equal_p (OMP_CLAUSE_DECL (c1),
+ OMP_CLAUSE_DECL (c2)))
+ && ((/* The clause size is generally not known right after
+ parsing. */
+ 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, remove duplicate map clauses
+ with identical kind only when the bits above
+ GOMP_MAP_FLAG_SPECIAL_2 are unset - as clauses with those flags
+ set 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 gcc/omp-general.h gcc/omp-general.h
index 4a6d58471de..a5126269650 100644
--- gcc/omp-general.h
+++ 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 gcc/omp-low.cc gcc/omp-low.cc
index b93012107f1..e8716d9fa37 100644
--- gcc/omp-low.cc
+++ 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 also 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 gcc/testsuite/c-c++-common/gomp/clauses-2.c gcc/testsuite/c-c++-common/gomp/clauses-2.c
index 8f98d57a312..d5018fa1ffa 100644
--- gcc/testsuite/c-c++-common/gomp/clauses-2.c
+++ 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 gcc/testsuite/c-c++-common/gomp/map-6.c gcc/testsuite/c-c++-common/gomp/map-6.c
index c1f3d1079b4..aa3e9786e5c 100644
--- gcc/testsuite/c-c++-common/gomp/map-6.c
+++ 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 gcc/testsuite/c-c++-common/gomp/map-multi-1.c gcc/testsuite/c-c++-common/gomp/map-multi-1.c
new file mode 100644
index 00000000000..496dcf6c77c
--- /dev/null
+++ 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 gcc/testsuite/c-c++-common/gomp/map-multi-2.c gcc/testsuite/c-c++-common/gomp/map-multi-2.c
new file mode 100644
index 00000000000..3eb4c1c004a
--- /dev/null
+++ 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 gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90 gcc/testsuite/gfortran.dg/gomp/map-multi-1.f90
new file mode 100644
index 00000000000..c12b6d431cf
--- /dev/null
+++ 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 gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90 gcc/testsuite/gfortran.dg/gomp/map-multi-2.f90
new file mode 100644
index 00000000000..fbd4dcbb256
--- /dev/null
+++ 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 gcc/testsuite/gfortran.dg/gomp/pr107214.f90 gcc/testsuite/gfortran.dg/gomp/pr107214.f90
index 25949934e84..84141b43727 100644
--- gcc/testsuite/gfortran.dg/gomp/pr107214.f90
+++ 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 libgomp/testsuite/libgomp.c++/target-map-multi-1.C libgomp/testsuite/libgomp.c++/target-map-multi-1.C
new file mode 100644
index 00000000000..54507176cf0
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c libgomp/testsuite/libgomp.c-c++-common/target-map-iterators-6.c
new file mode 100644
index 00000000000..b6013a99ac9
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c libgomp/testsuite/libgomp.c-c++-common/target-map-multi-1.c
new file mode 100644
index 00000000000..6d4d4c31f72
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c libgomp/testsuite/libgomp.c-c++-common/target-map-multi-2.c
new file mode 100644
index 00000000000..bb72dfcab49
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c libgomp/testsuite/libgomp.c-c++-common/target-map-multi-3.c
new file mode 100644
index 00000000000..c7aac627382
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c libgomp/testsuite/libgomp.c-c++-common/target-map-multi-4.c
new file mode 100644
index 00000000000..080ca93a88d
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90 libgomp/testsuite/libgomp.fortran/target-map-multi-1.f90
new file mode 100644
index 00000000000..0dc41bb136a
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90 libgomp/testsuite/libgomp.fortran/target-map-multi-2.f90
new file mode 100644
index 00000000000..783a7aa1f58
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90 libgomp/testsuite/libgomp.fortran/target-map-multi-3.f90
new file mode 100644
index 00000000000..44b601853d9
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90 libgomp/testsuite/libgomp.fortran/target-map-multi-4.f90
new file mode 100644
index 00000000000..88a14913611
--- /dev/null
+++ 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 libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90 libgomp/testsuite/libgomp.fortran/target-map-multi-5.f90
new file mode 100644
index 00000000000..94c7c941ed2
--- /dev/null
+++ 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
