head = NULL;
if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
false, NULL, &head,
- true) == MATCH_YES)
+ true, true) == MATCH_YES)
{
gfc_omp_namelist *n;
for (n = *head; n; n = n->next)
/* Look through component refs to find last array
reference. */
- if (openacc && resolved)
+ if (resolved)
{
/* The "!$acc cache" directive allows rectangular
subarrays to be specified, with some restrictions
arr(-n:n,-n:n) could be contiguous even if it looks
like it may not be. */
if (list != OMP_LIST_CACHE
+ && list != OMP_LIST_DEPEND
&& !gfc_is_simply_contiguous (n->expr, false, true)
&& gfc_is_not_contiguous (n->expr))
gfc_error ("Array is not contiguous at %L",
static void
gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
tree decl, bool element, gomp_map_kind ptr_kind,
- tree node, tree &node2, tree &node3, tree &node4)
+ tree &node, tree &node2, tree &node3, tree &node4)
{
gfc_se se;
tree ptr, ptr2;
+ tree elemsz = NULL_TREE;
gfc_init_se (&se, NULL);
gfc_conv_expr_reference (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
ptr = se.expr;
- OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+ OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
+ elemsz = OMP_CLAUSE_SIZE (node);
}
else
{
gfc_add_block_to_block (block, &se.pre);
OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
GFC_TYPE_ARRAY_RANK (type));
- tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
elemsz = fold_convert (gfc_array_index_type, elemsz);
OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type,
OMP_CLAUSE_SIZE (node), elemsz);
}
- gfc_add_block_to_block (block, &se.post);
+ gcc_assert (se.post.head == NULL_TREE);
ptr = fold_convert (build_pointer_type (char_type_node), ptr);
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+ ptr = fold_convert (ptrdiff_type_node, ptr);
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
OMP_CLAUSE_SIZE (node4) = size_int (0);
decl = build_fold_indirect_ref (decl);
}
- ptr = fold_convert (sizetype, ptr);
+ else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER
+ && n->expr->ts.type == BT_CHARACTER
+ && n->expr->ts.deferred)
+ {
+ gomp_map_kind map_kind;
+ if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
+ map_kind = GOMP_MAP_TO;
+ else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
+ map_kind = OMP_CLAUSE_MAP_KIND (node);
+ else
+ map_kind = GOMP_MAP_ALLOC;
+ gcc_assert (se.string_length);
+ node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
+ OMP_CLAUSE_DECL (node4) = se.string_length;
+ OMP_CLAUSE_SIZE (node4) = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+ }
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
+ tree desc_node;
tree type = TREE_TYPE (decl);
ptr2 = gfc_conv_descriptor_data_get (decl);
- node2 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
- OMP_CLAUSE_DECL (node2) = decl;
- OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_DECL (desc_node) = decl;
+ OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
+ if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
+ node2 = node;
+ node = desc_node; /* Needs to come first. */
+ }
+ else
+ {
+ OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
+ node2 = desc_node;
+ }
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
+ /* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra
+ cast prevents gimplify.c from recognising it as being part of the
+ struct – and adding an 'alloc: for the 'desc.data' pointer, which
+ would break as the 'desc' (the descriptor) is also mapped
+ (see node4 above). */
if (ptr_kind == GOMP_MAP_ATTACH_DETACH)
STRIP_NOPS (OMP_CLAUSE_DECL (node3));
}
else
{
if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- ptr2 = build_fold_addr_expr (decl);
+ {
+ tree offset;
+ ptr2 = build_fold_addr_expr (decl);
+ offset = fold_build2 (MINUS_EXPR, ptrdiff_type_node, ptr,
+ fold_convert (ptrdiff_type_node, ptr2));
+ offset = build2 (TRUNC_DIV_EXPR, ptrdiff_type_node,
+ offset, fold_convert (ptrdiff_type_node, elemsz));
+ offset = build4_loc (input_location, ARRAY_REF,
+ TREE_TYPE (TREE_TYPE (decl)),
+ decl, offset, NULL_TREE, NULL_TREE);
+ OMP_CLAUSE_DECL (node) = offset;
+ }
else
{
gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
OMP_CLAUSE_DECL (node3) = decl;
}
- ptr2 = fold_convert (sizetype, ptr2);
- OMP_CLAUSE_SIZE (node3)
- = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ ptr2 = fold_convert (ptrdiff_type_node, ptr2);
+ OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
+ ptr, ptr2);
}
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
- locus where, bool declare_simd = false)
+ locus where, bool declare_simd = false,
+ bool openacc = false)
{
tree omp_clauses = NULL_TREE, chunk_size, c;
int list, ifc;
tree node2 = NULL_TREE;
tree node3 = NULL_TREE;
tree node4 = NULL_TREE;
+
+ switch (n->u.map_op)
+ {
+ case OMP_MAP_ALLOC:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+ break;
+ case OMP_MAP_IF_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
+ break;
+ case OMP_MAP_ATTACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
+ break;
+ case OMP_MAP_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
+ break;
+ case OMP_MAP_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
+ break;
+ case OMP_MAP_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
+ break;
+ case OMP_MAP_ALWAYS_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
+ break;
+ case OMP_MAP_ALWAYS_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
+ break;
+ case OMP_MAP_ALWAYS_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
+ break;
+ case OMP_MAP_RELEASE:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
+ break;
+ case OMP_MAP_DELETE:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
+ break;
+ case OMP_MAP_DETACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
+ break;
+ case OMP_MAP_FORCE_ALLOC:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
+ break;
+ case OMP_MAP_FORCE_TO:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
+ break;
+ case OMP_MAP_FORCE_FROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
+ break;
+ case OMP_MAP_FORCE_TOFROM:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
+ break;
+ case OMP_MAP_FORCE_PRESENT:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
+ break;
+ case OMP_MAP_FORCE_DEVICEPTR:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
+ break;
+ default:
+ gcc_unreachable ();
+ }
+
tree decl = gfc_trans_omp_variable (n->sym, false);
if (DECL_P (decl))
TREE_ADDRESSABLE (decl) = 1;
&& n->expr->ref->u.ar.type == AR_FULL))
{
tree present = gfc_omp_check_optional_argument (decl, true);
- if (n->sym->ts.type == BT_CLASS)
+ if (openacc && n->sym->ts.type == BT_CLASS)
{
tree type = TREE_TYPE (decl);
if (n->sym->attr.optional)
/* Last component is a scalar. */
gfc_conv_expr (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
- OMP_CLAUSE_DECL (node) = se.expr;
+ /* For BT_CHARACTER a pointer is returned. */
+ OMP_CLAUSE_DECL (node)
+ = POINTER_TYPE_P (TREE_TYPE (se.expr))
+ ? build_fold_indirect_ref (se.expr) : se.expr;
gfc_add_block_to_block (block, &se.post);
+ if (sym_attr.pointer || sym_attr.allocatable)
+ {
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2,
+ openacc
+ ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER);
+ OMP_CLAUSE_DECL (node2)
+ = POINTER_TYPE_P (TREE_TYPE (se.expr))
+ ? se.expr : gfc_build_addr_expr (NULL, se.expr);
+ OMP_CLAUSE_SIZE (node2) = size_int (0);
+ if (!openacc
+ && n->expr->ts.type == BT_CHARACTER
+ && n->expr->ts.deferred)
+ {
+ gcc_assert (se.string_length);
+ tree tmp = gfc_get_char_type (n->expr->ts.kind);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, size_type_node,
+ fold_convert (size_type_node,
+ se.string_length),
+ TYPE_SIZE_UNIT (tmp));
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO);
+ OMP_CLAUSE_DECL (node3) = se.string_length;
+ OMP_CLAUSE_SIZE (node3)
+ = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+ }
+ }
goto finalize_map_clause;
}
if (lastcomp->u.c.component->ts.type == BT_DERIVED
|| lastcomp->u.c.component->ts.type == BT_CLASS)
{
- if (sym_attr.allocatable || sym_attr.pointer)
+ if (sym_attr.pointer || (openacc && sym_attr.allocatable))
{
tree data, size;
node2 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node2,
- GOMP_MAP_ATTACH_DETACH);
+ openacc
+ ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_DECL (node2) = data;
OMP_CLAUSE_SIZE (node2) = size_int (0);
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
{
+ gomp_map_kind map_kind;
+ tree desc_node;
tree type = TREE_TYPE (inner);
tree ptr = gfc_conv_descriptor_data_get (inner);
ptr = build_fold_indirect_ref (ptr);
OMP_CLAUSE_DECL (node) = ptr;
- node2 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
- OMP_CLAUSE_DECL (node2) = inner;
- OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3,
- GOMP_MAP_ATTACH_DETACH);
- OMP_CLAUSE_DECL (node3)
- = gfc_conv_descriptor_data_get (inner);
- STRIP_NOPS (OMP_CLAUSE_DECL (node3));
- OMP_CLAUSE_SIZE (node3) = size_int (0);
int rank = GFC_TYPE_ARRAY_RANK (type);
OMP_CLAUSE_SIZE (node)
= gfc_full_array_size (block, inner, rank);
tree elemsz
= TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
+ map_kind = GOMP_MAP_TO;
+ else if (n->u.map_op == OMP_MAP_RELEASE
+ || n->u.map_op == OMP_MAP_DELETE)
+ map_kind = OMP_CLAUSE_MAP_KIND (node);
+ else
+ map_kind = GOMP_MAP_ALLOC;
+ if (!openacc
+ && n->expr->ts.type == BT_CHARACTER
+ && n->expr->ts.deferred)
+ {
+ gcc_assert (se.string_length);
+ tree len = fold_convert (size_type_node,
+ se.string_length);
+ elemsz = gfc_get_char_type (n->expr->ts.kind);
+ elemsz = TYPE_SIZE_UNIT (elemsz);
+ elemsz = fold_build2 (MULT_EXPR, size_type_node,
+ len, elemsz);
+ node4 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
+ OMP_CLAUSE_DECL (node4) = se.string_length;
+ OMP_CLAUSE_SIZE (node4)
+ = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+ }
elemsz = fold_convert (gfc_array_index_type, elemsz);
OMP_CLAUSE_SIZE (node)
= fold_build2 (MULT_EXPR, gfc_array_index_type,
OMP_CLAUSE_SIZE (node), elemsz);
+ desc_node = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ if (openacc)
+ OMP_CLAUSE_SET_MAP_KIND (desc_node,
+ GOMP_MAP_TO_PSET);
+ else
+ OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
+ OMP_CLAUSE_DECL (desc_node) = inner;
+ OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
+ if (openacc)
+ node2 = desc_node;
+ else
+ {
+ node2 = node;
+ node = desc_node; /* Put first. */
+ }
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node3,
+ openacc
+ ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER);
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (inner);
+ /* Similar to gfc_trans_omp_array_section (details
+ there), we add/keep the cast for OpenMP to prevent
+ that an 'alloc:' gets added for node3 ('desc.data')
+ as that is part of the whole descriptor (node3).
+ TODO: Remove once the ME handles this properly. */
+ if (!openacc)
+ OMP_CLAUSE_DECL (node3)
+ = fold_convert (TREE_TYPE (TREE_OPERAND(ptr, 0)),
+ OMP_CLAUSE_DECL (node3));
+ else
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+ OMP_CLAUSE_SIZE (node3) = size_int (0);
}
else
OMP_CLAUSE_DECL (node) = inner;
&& lastcomp->next->type == REF_ARRAY
&& lastcomp->next->u.ar.type == AR_ELEMENT);
+ gomp_map_kind kind = (openacc ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER);
gfc_trans_omp_array_section (block, n, inner, element,
- GOMP_MAP_ATTACH_DETACH,
- node, node2, node3, node4);
+ kind, node, node2, node3,
+ node4);
}
}
else /* An array element or array section. */
}
finalize_map_clause:
- switch (n->u.map_op)
- {
- case OMP_MAP_ALLOC:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
- break;
- case OMP_MAP_IF_PRESENT:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
- break;
- case OMP_MAP_ATTACH:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
- break;
- case OMP_MAP_TO:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
- break;
- case OMP_MAP_FROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
- break;
- case OMP_MAP_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
- break;
- case OMP_MAP_ALWAYS_TO:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
- break;
- case OMP_MAP_ALWAYS_FROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
- break;
- case OMP_MAP_ALWAYS_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
- break;
- case OMP_MAP_RELEASE:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
- break;
- case OMP_MAP_DELETE:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
- break;
- case OMP_MAP_DETACH:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
- break;
- case OMP_MAP_FORCE_ALLOC:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
- break;
- case OMP_MAP_FORCE_TO:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
- break;
- case OMP_MAP_FORCE_FROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
- break;
- case OMP_MAP_FORCE_TOFROM:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
- break;
- case OMP_MAP_FORCE_PRESENT:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
- break;
- case OMP_MAP_FORCE_DEVICEPTR:
- OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
- break;
- default:
- gcc_unreachable ();
- }
+
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
if (node2)
omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
- code->loc);
+ code->loc, false, true);
stmt = gfc_trans_omp_code (code->block->next, true);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
- code->loc);
+ code->loc, false, true);
stmt = build1_loc (input_location, construct_code, void_type_node,
oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
if (construct_code == OACC_KERNELS)
construct_clauses.lists[OMP_LIST_REDUCTION] = NULL;
oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
- code->loc);
+ code->loc, false, true);
}
if (!loop_clauses.seq)
pblock = █
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
- code->loc);
+ code->loc, false, true);
stmt = gfc_trans_omp_code (code->block->next, true);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm\\.0\\.data - \\(integer\\(kind=8\\)\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
-! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm\\.1\\.data - \\(integer\\(kind=8\\)\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f
!$omp target map(j(:))
!$omp end target
- !$omp target map(j(1:9:2)) ! { dg-error "Stride should not be specified for array section in MAP clause" }
+ !$omp target map(j(1:9:2))
+ ! { dg-error "Array is not contiguous" "" { target *-*-* } 60 }
+ ! { dg-error "Stride should not be specified for array section in MAP clause" "" { target *-*-* } 60 }
!$omp end target
!$omp target map(aas(5:))
!$omp end target
- ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 63 }
- ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 63 }
+ ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 65 }
+ ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 65 }
!$omp target map(aas(:))
!$omp end target
- ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 68 }
- ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 68 }
+ ! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 70 }
+ ! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 70 }
!$omp target map(aas) ! { dg-error "Assumed size array" }
!$omp end target
!$omp target map(k(5:))
!$omp end target
- ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 82 }
- ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 82 }
+ ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 84 }
+ ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 84 }
!$omp target map(k(5:,:,3))
!$omp end target
- ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 87 }
- ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 87 }
+ ! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 89 }
+ ! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 89 }
!$omp target map(tt)
!$omp end target
- !$omp target map(tt%i) ! { dg-error "Syntax error in OpenMP variable list" }
+ !$omp target map(tt%k) ! { dg-error "not a member of" }
!$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
- !$omp target map(tt%j) ! { dg-error "Syntax error in OpenMP variable list" }
- !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+ !$omp target map(tt%j)
+ !$omp end target
- ! broken test
- !$omp target map(tt%j(1)) ! { dg-error "Syntax error in OpenMP variable list" }
- !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+ !$omp target map(tt%j(1))
+ !$omp end target
- !$omp target map(tt%j(1:)) ! { dg-error "Syntax error in OpenMP variable list" }
- !$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
+ !$omp target map(tt%j(1:))
+ !$omp end target
!$omp target map(tp) ! { dg-error "THREADPRIVATE object 'tp' in MAP clause" }
!$omp end target
--- /dev/null
+type t
+ integer :: i
+end type t
+type(t) v
+!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" }
+end
--- /dev/null
+! { dg-do run }
+!
+! Test OpenMP 4.5 structure-element mapping
+
+! TODO: character(kind=4,...) needs to be tested, but depends on
+! PR fortran/95837
+! TODO: ...%str4 should be tested but that currently fails due to
+! PR fortran/95868 (see commented lined)
+! TODO: Test also array-valued var, nested derived types,
+! type-extended types.
+
+program main
+ implicit none
+
+ type t2
+ integer :: a, b
+ ! For complex, assume small integers are exactly representable
+ complex(kind=8) :: c
+ integer :: d(10)
+ integer, pointer :: e => null(), f(:) => null()
+ character(len=5) :: str1
+ character(len=5) :: str2(4)
+ character(len=:), pointer :: str3 => null()
+ character(len=:), pointer :: str4(:) => null()
+ end type t2
+
+ integer :: i
+
+ call one ()
+ call two ()
+ call three ()
+ call four ()
+ call five ()
+ call six ()
+ call seven ()
+ call eight ()
+
+contains
+ ! Implicitly mapped – but no pointers are mapped
+ subroutine one()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "one" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%e, source=99)
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str3, source="HelloWorld")
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom:var)
+ if (var%a /= 1) stop 1
+ if (var%b /= 2) stop 2
+ if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+ if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+ if (var%str1 /= "abcde") stop 5
+ if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+ !$omp end target
+
+ deallocate(var%e, var%f, var%str3, var%str4)
+ end subroutine one
+
+ ! Explicitly mapped – all and full arrays
+ subroutine two()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "two" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%e, source=99)
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str3, source="HelloWorld")
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%a, var%b, var%c, var%d, var%e, var%f, &
+ !$omp& var%str1, var%str2, var%str3, var%str4)
+ if (var%a /= 1) stop 1
+ if (var%b /= 2) stop 2
+ if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+ if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+ if (var%str1 /= "abcde") stop 5
+ if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+
+ if (.not. associated (var%e)) stop 7
+ if (var%e /= 99) stop 8
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f /= [22, 33, 44, 55])) stop 11
+ if (.not. associated (var%str3)) stop 12
+ if (len (var%str3) /= len ("HelloWorld")) stop 13
+ if (var%str3 /= "HelloWorld") stop 14
+ if (.not. associated (var%str4)) stop 15
+ if (len (var%str4) /= 5) stop 16
+ if (size (var%str4) /= 2) stop 17
+ if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
+ !$omp end target
+
+ deallocate(var%e, var%f, var%str3, var%str4)
+ end subroutine two
+
+ ! Explicitly mapped – one by one but full arrays
+ subroutine three()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "three" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%e, source=99)
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str3, source="HelloWorld")
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%a)
+ if (var%a /= 1) stop 1
+ !$omp end target
+ !$omp target map(tofrom: var%b)
+ if (var%b /= 2) stop 2
+ !$omp end target
+ !$omp target map(tofrom: var%c)
+ if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
+ !$omp end target
+ !$omp target map(tofrom: var%d)
+ if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str1)
+ if (var%str1 /= "abcde") stop 5
+ !$omp end target
+ !$omp target map(tofrom: var%str2)
+ if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
+ !$omp end target
+
+ !$omp target map(tofrom: var%e)
+ if (.not. associated (var%e)) stop 7
+ if (var%e /= 99) stop 8
+ !$omp end target
+ !$omp target map(tofrom: var%f)
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f /= [22, 33, 44, 55])) stop 11
+ !$omp end target
+ !$omp target map(tofrom: var%str3)
+ if (.not. associated (var%str3)) stop 12
+ if (len (var%str3) /= len ("HelloWorld")) stop 13
+ if (var%str3 /= "HelloWorld") stop 14
+ !$omp end target
+ !$omp target map(tofrom: var%str4)
+ if (.not. associated (var%str4)) stop 15
+ if (len (var%str4) /= 5) stop 16
+ if (size (var%str4) /= 2) stop 17
+ if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
+ !$omp end target
+
+ deallocate(var%e, var%f, var%str3, var%str4)
+ end subroutine three
+
+ ! Explicitly mapped – all but only subarrays
+ subroutine four()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "four" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3), var%str4(2:2))
+ !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+ !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine four
+
+ ! Explicitly mapped – all but only subarrays and one by one
+ subroutine five()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "five" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%d(4:7))
+ if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+
+ !$omp target map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine five
+
+ ! Explicitly mapped – all but only array elements
+ subroutine six()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "six" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), var%str4(2))
+ !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3))
+ if (var%d(5) /= -3*5) stop 4
+ if (var%str2(3) /= "ABCDE") stop 6
+
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (var%f(3) /= 44) stop 11
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+ !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine six
+
+ ! Explicitly mapped – all but only array elements and one by one
+ subroutine seven()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "seven" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+ allocate (var%f, source=[22, 33, 44, 55])
+ allocate (var%str4, source=["Let's", "Go!!!"])
+
+ !$omp target map(tofrom: var%d(5))
+ if (var%d(5) /= (-3*5)) stop 4
+ !$omp end target
+ !$omp target map(tofrom: var%str2(2:3))
+ if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
+ !$omp end target
+
+ !$omp target map(tofrom: var%f(2:3))
+ if (.not. associated (var%f)) stop 9
+ if (size (var%f) /= 4) stop 10
+ if (any (var%f(2:3) /= [33, 44])) stop 11
+ !$omp end target
+! !$omp target map(tofrom: var%str4(2:2))
+! if (.not. associated (var%str4)) stop 15
+! if (len (var%str4) /= 5) stop 16
+! if (size (var%str4) /= 2) stop 17
+! if (var%str4(2) /= "Go!!!") stop 18
+! !$omp end target
+
+ deallocate(var%f, var%str4)
+ end subroutine seven
+
+ ! Check mapping of NULL pointers
+ subroutine eight()
+ type(t2) :: var, var2(4)
+ type(t2), pointer :: var3, var4(:)
+
+ print '(g0)', '==== TESTCASE "eight" ===='
+
+ var = t2(a = 1, &
+ b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
+ d = [(-3*i, i = 1, 10)], &
+ str1 = "abcde", &
+ str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
+
+! !$omp target map(tofrom: var%e, var%f, var%str3, var%str4)
+ !$omp target map(tofrom: var%e, var%str3)
+ if (associated (var%e)) stop 1
+! if (associated (var%f)) stop 2
+ if (associated (var%str3)) stop 3
+! if (associated (var%str4)) stop 4
+ !$omp end target
+ end subroutine eight
+
+end program main
projects / gcc.git / commitdiff