tree init_no_targetsync_clause = NULL_TREE;
tree depend_clause = NULL_TREE;
+ if (!openacc)
+ clauses = omp_remove_duplicate_maps (clauses, true);
+
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
bitmap_initialize (&firstprivate_head, &bitmap_default_obstack);
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)
tree init_no_targetsync_clause = NULL_TREE;
tree depend_clause = NULL_TREE;
+ if (!openacc)
+ clauses = omp_remove_duplicate_maps (clauses, true);
+
bitmap_obstack_initialize (NULL);
bitmap_initialize (&generic_head, &bitmap_default_obstack);
bitmap_initialize (&firstprivate_head, &bitmap_default_obstack);
{
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)))
+ || (openacc
+ && bitmap_bit_p (&map_firstprivate_head, DECL_UID (t))))
{
error_at (OMP_CLAUSE_LOCATION (c),
"%qD appears more than once in data clauses", 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_TARGET
&& ort != C_ORT_OMP_EXIT_DATA)
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
}
return false;
+ case OMP_ARRAY_SECTION:
+ return OP_SAME (0) && OP_SAME_WITH_NULL (1) && OP_SAME_WITH_NULL (2);
+
default:
return false;
}
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
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;
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 ((openacc || list != OMP_LIST_MAP) && n->sym->mark)
gfc_error ("Symbol %qs present on multiple clauses at %L",
n->sym->name, &n->where);
else
case OMP_LIST_MAP:
for (; n != NULL; n = n->next)
{
+ if (!openacc)
+ {
+ // 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;
|| code == OACC_UPDATE
|| code == OACC_DECLARE)
{
+ if (!(region_type & ORT_ACC))
+ *list_p = omp_remove_duplicate_maps (*list_p, false);
groups = omp_gather_mapping_groups (list_p);
if (groups)
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;
+}
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 */
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)
- {
- key = (splay_tree_key) &DECL_NAME (var);
- gcc_checking_assert (key != (splay_tree_key) var);
- }
- if ((mask & 8) != 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_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));
|| (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)
{
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);
}
&& 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)))
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);
}
}
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;
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);
continue;
}
- if (!maybe_lookup_field (var, ctx))
+ if (!maybe_lookup_field (c, ctx))
continue;
/* Don't remap compute constructs' reduction variables, because the
&& 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
}
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
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)))
}
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
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)))
{
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);
#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)
--- /dev/null
+/* { 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" } } */
--- /dev/null
+/* { 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" } } */
--- /dev/null
+! { 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" } }
--- /dev/null
+! { 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" } }
+
+
! { 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" } }
--- /dev/null
+// 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;
+}
--- /dev/null
+/* { 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);
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* Test multiple map clauses for the same variable with the 'always'
+ map-type modifier. */
+/* { dg-do run { target offload_device_nonshared_as } } */
+
+#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;
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+/* 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;
+}
--- /dev/null
+! 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
--- /dev/null
+! Test multiple map clauses for the same variable with the 'always'
+! map-type modifier.
+! { dg-do run { target offload_device_nonshared_as } }
+
+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
--- /dev/null
+! 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
--- /dev/null
+! 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
--- /dev/null
+! 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
projects / thirdparty / gcc.git / commitdiff
