/* Add attributes to variables. Functions are handled elsewhere. */
attributes = add_attributes_to_decl (sym->attr, NULL_TREE);
decl_attributes (&decl, attributes, 0);
+ if (sym->ts.deferred && VAR_P (length))
+ decl_attributes (&length, attributes, 0);
/* Symbols from modules should have their assembler names mangled.
This is done here rather than in gfc_finish_var_decl because it
/* Translate an array section or array element. */
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)
+gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op,
+ gfc_omp_namelist *n, tree decl, bool element,
+ gomp_map_kind ptr_kind, tree &node, tree &node2,
+ tree &node3, tree &node4)
{
gfc_se se;
tree ptr, ptr2;
tree elemsz = NULL_TREE;
gfc_init_se (&se, NULL);
-
if (element)
{
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 (TREE_TYPE (ptr)));
- elemsz = OMP_CLAUSE_SIZE (node);
}
else
{
gfc_conv_expr_descriptor (&se, n->expr);
ptr = gfc_conv_array_data (se.expr);
+ }
+ if (n->expr->ts.type == BT_CHARACTER && n->expr->ts.deferred)
+ {
+ gcc_assert (se.string_length);
+ tree len = gfc_evaluate_now (se.string_length, block);
+ elemsz = gfc_get_char_type (n->expr->ts.kind);
+ elemsz = TYPE_SIZE_UNIT (elemsz);
+ elemsz = fold_build2 (MULT_EXPR, size_type_node,
+ fold_convert (size_type_node, len), elemsz);
+ }
+ if (element)
+ {
+ if (!elemsz)
+ elemsz = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
+ OMP_CLAUSE_SIZE (node) = elemsz;
+ }
+ else
+ {
tree type = TREE_TYPE (se.expr);
gfc_add_block_to_block (block, &se.pre);
OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
GFC_TYPE_ARRAY_RANK (type));
- elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ if (!elemsz)
+ 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);
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
- && ptr_kind == GOMP_MAP_POINTER)
+ && ptr_kind == GOMP_MAP_POINTER
+ && op != EXEC_OMP_TARGET_EXIT_DATA
+ && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_RELEASE
+ && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_DELETE)
+
{
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
&& 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)
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
map_kind = OMP_CLAUSE_MAP_KIND (node);
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA
+ || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE)
+ map_kind = GOMP_MAP_RELEASE;
else
- map_kind = GOMP_MAP_ALLOC;
+ map_kind = GOMP_MAP_TO;
gcc_assert (se.string_length);
node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
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)
+ if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_DELETE);
+ node2 = desc_node;
+ }
+ else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
+ || op == EXEC_OMP_TARGET_EXIT_DATA)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_RELEASE);
+ node2 = desc_node;
+ }
+ else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
{
OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
node2 = node;
OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
node2 = desc_node;
}
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
+ if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ return;
+ 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);
+ OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl);
/* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra
cast prevents gimplify.cc from recognising it as being part of the
struct - and adding an 'alloc: for the 'desc.data' pointer, which
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
- bool openacc = false)
+ bool openacc = false, gfc_exec_op op = EXEC_NOP)
{
tree omp_clauses = NULL_TREE, prev_clauses, chunk_size, c;
tree iterator = NULL_TREE;
tree node2 = NULL_TREE;
tree node3 = NULL_TREE;
tree node4 = NULL_TREE;
+ tree node5 = NULL_TREE;
/* OpenMP: automatically map pointer targets with the pointer;
hence, always update the descriptor/pointer itself. */
|| (n->expr->ref->type == REF_ARRAY
&& n->expr->ref->u.ar.type == AR_FULL))
{
+ gomp_map_kind map_kind;
+ tree type = TREE_TYPE (decl);
+ if (n->sym->ts.type == BT_CHARACTER
+ && n->sym->ts.deferred
+ && n->sym->attr.omp_declare_target
+ && (always_modifier || n->sym->attr.pointer)
+ && op != EXEC_OMP_TARGET_EXIT_DATA
+ && n->u.map_op != OMP_MAP_DELETE
+ && n->u.map_op != OMP_MAP_RELEASE)
+ {
+ gcc_assert (n->sym->ts.u.cl->backend_decl);
+ node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node5, GOMP_MAP_ALWAYS_TO);
+ OMP_CLAUSE_DECL (node5) = n->sym->ts.u.cl->backend_decl;
+ OMP_CLAUSE_SIZE (node5)
+ = TYPE_SIZE_UNIT (gfc_charlen_type_node);
+ }
+
tree present = gfc_omp_check_optional_argument (decl, true);
if (openacc && n->sym->ts.type == BT_CLASS)
{
OMP_CLAUSE_SIZE (node2) = size_int (0);
goto finalize_map_clause;
}
- else if (POINTER_TYPE_P (TREE_TYPE (decl))
+ else if (POINTER_TYPE_P (type)
&& (gfc_omp_privatize_by_reference (decl)
|| GFC_DECL_GET_SCALAR_POINTER (decl)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
|| GFC_DECL_CRAY_POINTEE (decl)
- || GFC_DESCRIPTOR_TYPE_P
- (TREE_TYPE (TREE_TYPE (decl)))
+ || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))
|| (n->sym->ts.type == BT_DERIVED
&& (n->sym->ts.u.derived->ts.f90_type
!= BT_VOID))))
to avoid accessing undefined variables. We cannot use
a temporary variable here as otherwise the replacement
of the variables in omp-low.cc will not work. */
- if (present && GFC_ARRAY_TYPE_P (TREE_TYPE (decl)))
+ if (present && GFC_ARRAY_TYPE_P (type))
{
tree tmp = fold_build2_loc (input_location,
MODIFY_EXPR,
cond, tmp,
NULL_TREE));
}
- node4 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
- OMP_CLAUSE_DECL (node4) = decl;
- OMP_CLAUSE_SIZE (node4) = size_int (0);
+ /* For descriptor types, the unmapping happens below. */
+ if (op != EXEC_OMP_TARGET_EXIT_DATA
+ || !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ enum gomp_map_kind gmk = GOMP_MAP_POINTER;
+ if (op == EXEC_OMP_TARGET_EXIT_DATA
+ && n->u.map_op == OMP_MAP_DELETE)
+ gmk = GOMP_MAP_DELETE;
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ gmk = GOMP_MAP_RELEASE;
+ tree size;
+ if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+ size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ else
+ size = size_int (0);
+ node4 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node4, gmk);
+ OMP_CLAUSE_DECL (node4) = decl;
+ OMP_CLAUSE_SIZE (node4) = size;
+ }
decl = build_fold_indirect_ref (decl);
if ((TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE
|| gfc_omp_is_optional_argument (orig_decl))
&& (GFC_DECL_GET_SCALAR_POINTER (orig_decl)
|| GFC_DECL_GET_SCALAR_ALLOCATABLE (orig_decl)))
{
+ enum gomp_map_kind gmk;
+ if (op == EXEC_OMP_TARGET_EXIT_DATA
+ && n->u.map_op == OMP_MAP_DELETE)
+ gmk = GOMP_MAP_DELETE;
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ gmk = GOMP_MAP_RELEASE;
+ else
+ gmk = GOMP_MAP_POINTER;
+ tree size;
+ if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE)
+ size = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ else
+ size = size_int (0);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, gmk);
OMP_CLAUSE_DECL (node3) = decl;
- OMP_CLAUSE_SIZE (node3) = size_int (0);
+ OMP_CLAUSE_SIZE (node3) = size;
decl = build_fold_indirect_ref (decl);
}
}
gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr)));
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);
+ node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_DECL (node2) = decl;
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- if (present)
- {
- ptr = gfc_conv_descriptor_data_get (decl);
- ptr = gfc_build_addr_expr (NULL, ptr);
- ptr = gfc_build_cond_assign_expr (block, present, ptr,
- null_pointer_node);
- ptr = build_fold_indirect_ref (ptr);
- OMP_CLAUSE_DECL (node3) = ptr;
- }
+ if (n->u.map_op == OMP_MAP_DELETE)
+ map_kind = GOMP_MAP_DELETE;
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA
+ || n->u.map_op == OMP_MAP_RELEASE)
+ map_kind = GOMP_MAP_RELEASE;
else
- OMP_CLAUSE_DECL (node3)
- = gfc_conv_descriptor_data_get (decl);
- OMP_CLAUSE_SIZE (node3) = size_int (0);
- if (n->u.map_op == OMP_MAP_ATTACH)
- {
- /* Standalone attach clauses used with arrays with
- descriptors must copy the descriptor to the target,
- else they won't have anything to perform the
- attachment onto (see OpenACC 2.6, "2.6.3. Data
- Structures with Pointers"). */
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
- /* We don't want to map PTR at all in this case, so
- delete its node and shuffle the others down. */
- node = node2;
- node2 = node3;
- node3 = NULL;
- goto finalize_map_clause;
- }
- else if (n->u.map_op == OMP_MAP_DETACH)
+ map_kind = GOMP_MAP_TO_PSET;
+ OMP_CLAUSE_SET_MAP_KIND (node2, map_kind);
+
+ if (op != EXEC_OMP_TARGET_EXIT_DATA
+ && n->u.map_op != OMP_MAP_DELETE
+ && n->u.map_op != OMP_MAP_RELEASE)
{
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
- /* Similarly to above, we don't want to unmap PTR
- here. */
- node = node2;
- node2 = node3;
- node3 = NULL;
- goto finalize_map_clause;
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ if (present)
+ {
+ ptr = gfc_conv_descriptor_data_get (decl);
+ ptr = gfc_build_addr_expr (NULL, ptr);
+ ptr = gfc_build_cond_assign_expr (
+ block, present, ptr, null_pointer_node);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (node3) = ptr;
+ }
+ else
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (decl);
+ OMP_CLAUSE_SIZE (node3) = size_int (0);
+
+ if (n->u.map_op == OMP_MAP_ATTACH)
+ {
+ /* Standalone attach clauses used with arrays with
+ descriptors must copy the descriptor to the
+ target, else they won't have anything to
+ perform the attachment onto (see OpenACC 2.6,
+ "2.6.3. Data Structures with Pointers"). */
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH);
+ /* We don't want to map PTR at all in this case,
+ so delete its node and shuffle the others
+ down. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else if (n->u.map_op == OMP_MAP_DETACH)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH);
+ /* Similarly to above, we don't want to unmap PTR
+ here. */
+ node = node2;
+ node2 = node3;
+ node3 = NULL;
+ goto finalize_map_clause;
+ }
+ else
+ OMP_CLAUSE_SET_MAP_KIND (node3,
+ always_modifier
+ ? GOMP_MAP_ALWAYS_POINTER
+ : GOMP_MAP_POINTER);
}
- else
- OMP_CLAUSE_SET_MAP_KIND (node3,
- always_modifier
- ? GOMP_MAP_ALWAYS_POINTER
- : GOMP_MAP_POINTER);
/* We have to check for n->sym->attr.dimension because
of scalar coarrays. */
tem
= gfc_full_array_size (&cond_block, decl,
GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz;
+ if (n->sym->ts.type == BT_CHARACTER
+ && n->sym->ts.deferred)
+ {
+ tree len = n->sym->ts.u.cl->backend_decl;
+ len = fold_convert (size_type_node, len);
+ elemsz = gfc_get_char_type (n->sym->ts.kind);
+ elemsz = TYPE_SIZE_UNIT (elemsz);
+ elemsz = fold_build2 (MULT_EXPR, size_type_node,
+ len, elemsz);
+ }
+ else
+ elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ tem = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ tem, elemsz);
gfc_add_modify (&cond_block, size, tem);
then_b = gfc_finish_block (&cond_block);
gfc_init_block (&cond_block);
gfc_init_block (&cond_block);
tree size = gfc_full_array_size (&cond_block, decl,
GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ size = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ size, elemsz);
+ size = gfc_evaluate_now (size, &cond_block);
if (present)
{
tree var = gfc_create_var (gfc_array_index_type,
OMP_CLAUSE_SIZE (node) = size;
}
}
- if (n->sym->attr.dimension)
- {
- tree 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);
- }
}
else if (present
&& TREE_CODE (decl) == INDIRECT_REF
}
else
OMP_CLAUSE_DECL (node) = decl;
+
+ if (!n->sym->attr.dimension
+ && n->sym->ts.type == BT_CHARACTER
+ && n->sym->ts.deferred)
+ {
+ if (!DECL_P (decl))
+ {
+ gcc_assert (TREE_CODE (decl) == INDIRECT_REF);
+ decl = TREE_OPERAND (decl, 0);
+ }
+ tree cond = fold_build2_loc (input_location, NE_EXPR,
+ boolean_type_node,
+ decl, null_pointer_node);
+ if (present)
+ cond = fold_build2_loc (input_location,
+ TRUTH_ANDIF_EXPR,
+ boolean_type_node,
+ present, cond);
+ tree len = n->sym->ts.u.cl->backend_decl;
+ len = fold_convert (size_type_node, len);
+ tree size = gfc_get_char_type (n->sym->ts.kind);
+ size = TYPE_SIZE_UNIT (size);
+ size = fold_build2 (MULT_EXPR, size_type_node, len, size);
+ size = build3_loc (input_location,
+ COND_EXPR,
+ size_type_node,
+ cond, size,
+ size_zero_node);
+ size = gfc_evaluate_now (size, block);
+ OMP_CLAUSE_SIZE (node) = size;
+ }
}
else if (n->expr
&& n->expr->expr_type == EXPR_VARIABLE
&& !(POINTER_TYPE_P (type)
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type))))
k = GOMP_MAP_FIRSTPRIVATE_POINTER;
- gfc_trans_omp_array_section (block, n, decl, element, k,
+ gfc_trans_omp_array_section (block, op, n, decl, element, k,
node, node2, node3, node4);
}
else if (n->expr
fold_convert (size_type_node,
se.string_length),
TYPE_SIZE_UNIT (tmp));
+ if (n->u.map_op == OMP_MAP_DELETE)
+ kind = GOMP_MAP_DELETE;
+ else if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ kind = GOMP_MAP_RELEASE;
+ else
+ kind = GOMP_MAP_TO;
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO);
+ OMP_CLAUSE_SET_MAP_KIND (node3, kind);
OMP_CLAUSE_DECL (node3) = se.string_length;
OMP_CLAUSE_SIZE (node3)
= TYPE_SIZE_UNIT (gfc_charlen_type_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;
+ map_kind = OMP_CLAUSE_MAP_KIND (node);
+ if (GOMP_MAP_COPY_TO_P (map_kind)
+ || map_kind == GOMP_MAP_ALLOC)
+ map_kind = ((GOMP_MAP_ALWAYS_P (map_kind)
+ || gfc_expr_attr (n->expr).pointer)
+ ? GOMP_MAP_ALWAYS_TO : 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 if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ map_kind = GOMP_MAP_RELEASE;
else
map_kind = GOMP_MAP_ALLOC;
if (!openacc
node2 = node;
node = desc_node; /* Put first. */
}
+ if (op == EXEC_OMP_TARGET_EXIT_DATA)
+ goto finalize_map_clause;
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3,
bool element = lastref->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,
+ gfc_trans_omp_array_section (block, op, n, inner, element,
kind, node, node2, node3,
node4);
}
omp_clauses = gfc_trans_add_clause (node3, omp_clauses);
if (node4)
omp_clauses = gfc_trans_add_clause (node4, omp_clauses);
+ if (node5)
+ omp_clauses = gfc_trans_add_clause (node5, omp_clauses);
}
break;
case OMP_LIST_TO:
gfc_start_block (&block);
omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
- code->loc);
+ code->loc, false, false, code->op);
stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node,
omp_clauses);
gfc_add_expr_to_block (&block, stmt);
{
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
- tree *prev_list_p = NULL, *orig_list_p = list_p;
+ tree *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
default:
break;
}
- /* For Fortran, not only the pointer to the data is mapped but also
- the address of the pointer, the array descriptor etc.; for
- 'exit data' - and in particular for 'delete:' - having an 'alloc:'
- does not make sense. Likewise, for 'update' only transferring the
- data itself is needed as the rest has been handled in previous
- directives. However, for 'exit data', the array descriptor needs
- to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE.
-
- NOTE: Generally, it is not safe to perform "enter data" operations
- on arrays where the data *or the descriptor* may go out of scope
- before a corresponding "exit data" operation -- and such a
- descriptor may be synthesized temporarily, e.g. to pass an
- explicit-shape array to a function expecting an assumed-shape
- argument. Performing "enter data" inside the called function
- would thus be problematic. */
- if (code == OMP_TARGET_EXIT_DATA
- && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)
- OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p)
- == GOMP_MAP_DELETE
- ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE);
- else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE)
- && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET))
- remove = true;
-
if (remove)
break;
if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC))
remove = true;
break;
}
-
- if (!remove
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH
- && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
- && OMP_CLAUSE_CHAIN (c)
- && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
- && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ALWAYS_POINTER)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ATTACH_DETACH)
- || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_TO_PSET)))
- prev_list_p = list_p;
-
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) 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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
-! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
+! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[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 acc exit data map\\(from:cpo_r\\);$" 1 "original" } }
--- /dev/null
+! { dg-additional-options "-fdump-tree-omplower" }
+
+! If enter data adds a (GOMP_MAP_)POINTER attachment, exit data needs to remove
+! it again. If not there can be all kind of issues, in particular when
+! stack memory was mapped, reused later and mapped again.
+
+subroutine test_aa (aa2, aa3)
+ integer(kind=4), allocatable :: aa1, aa2, aa3
+ optional :: aa3
+ !$omp target enter data map(aa1)
+ !$omp target exit data map(aa1)
+ !$omp target enter data map(aa2)
+ !$omp target exit data map(aa2)
+ !$omp target enter data map(aa3)
+ !$omp target exit data map(aa3)
+end
+
+subroutine test_pp (pp2, pp3)
+ integer(kind=4), allocatable :: pp1, pp2, pp3
+ optional :: pp3
+ !$omp target enter data map(pp1)
+ !$omp target exit data map(pp1)
+ !$omp target enter data map(pp2)
+ !$omp target exit data map(pp2)
+ !$omp target enter data map(pp3)
+ !$omp target exit data map(pp3)
+end
+
+subroutine test_pprelease (rp2, rp3)
+ integer(kind=4), allocatable :: rp1, rp2, rp3
+ optional :: rp3
+ !$omp target enter data map(rp1)
+ !$omp target exit data map(release:rp1)
+ !$omp target enter data map(rp2)
+ !$omp target exit data map(release:rp2)
+ !$omp target enter data map(rp3)
+ !$omp target exit data map(release:rp3)
+end
+
+subroutine test_ppdelete (dp2, dp3)
+ integer(kind=4), allocatable :: dp1, dp2, dp3
+ optional :: dp3
+ !$omp target enter data map(dp1)
+ !$omp target exit data map(delete:dp1)
+ !$omp target enter data map(dp2)
+ !$omp target exit data map(delete:dp2)
+ !$omp target enter data map(dp3)
+ !$omp target exit data map(delete:dp3)
+end
+
+
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:aa1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa2 \\\[len: .\\\]\\) map\\(release:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa3 \\\[len: .\\\]\\) map\\(release:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:pp1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp1 \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp2 \\\[len: .\\\]\\) map\\(release:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp3 \\\[len: .\\\]\\) map\\(release:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp1 \\\[len: .\\\]\\) map\\(release:\\*rp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp2 \\\[len: .\\\]\\) map\\(release:\\*rp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:rp3 \\\[len: .\\\]\\) map\\(release:\\*rp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp1 \\\[len: .\\\]\\) map\\(delete:\\*dp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp2 \\\[len: .\\\]\\) map\\(delete:\\*dp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(delete:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:dp3 \\\[len: .\\\]\\) map\\(delete:\\*dp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } }
! PR fortran/108545
-! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM <integer\\(kind=4\\)\\\[0:\\\]> \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } }
program p
type t
end subroutine sub
end module m
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! Check for multiplication: len = arrays_size * 4<bytes>:
+! { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = D\\.\[0-9\]+ \\* 4;" 2 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
-! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
+! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } }
! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } }
!$omp target exit data map(from:three)
end
-! { dg-final { scan-tree-dump "omp target exit data map\\(delete:.*\\) map\\(delete:one \\\[len: .*\\\]\\)" "omplower" } }
-! { dg-final { scan-tree-dump "omp target exit data map\\(release:.*\\) map\\(release:two \\\[len: .*\\\]\\)" "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:one \\\[len: \[0-9\]+\\\]\\) map\\(delete:MEM " "omplower" } }
+! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:two \\\[len: \[0-9\]+\\\]\\) map\\(release:MEM " "omplower" } }
! { dg-final { scan-tree-dump "omp target exit data map\\(from:.*\\) map\\(release:three \\\[len: .*\\\]\\)" "omplower" } }
if (any (var%p1 /= [22,53,28,6,4])) stop 3
if (any (var%p2 /= [46,679,54])) stop 4
!$omp end target
-!!$omp target exit data map(from:var%p1, var%p2)
+!$omp target exit data map(from:var%p1, var%p2)
end
--- /dev/null
+! Check that 'map(alloc:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+ integer :: ic(2:5), ic2
+ character(len=11) :: ccstr(3:4), ccstr2
+ character(len=11,kind=4) :: cc4str(3:7), cc4str2
+ integer, pointer :: pc(:), pc2
+ character(len=:), pointer :: pcstr(:), pcstr2
+ character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+end type t
+
+type(t) :: dt
+
+integer :: ii(5), ii2
+character(len=11) :: clstr(-1:1), clstr2
+character(len=11,kind=4) :: cl4str(0:3), cl4str2
+integer, pointer :: ip(:), ip2
+integer, allocatable :: ia(:), ia2
+character(len=:), pointer :: pstr(:), pstr2
+character(len=:), allocatable :: astr(:), astr2
+character(len=:,kind=4), pointer :: p4str(:), p4str2
+character(len=:,kind=4), allocatable :: a4str(:), a4str2
+
+allocate(dt%pc(5), dt%pc2)
+allocate(character(len=2) :: dt%pcstr(2))
+allocate(character(len=4) :: dt%pcstr2)
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+allocate(ip(5), ip2, ia(8), ia2)
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=4) :: pstr2)
+allocate(character(len=6) :: astr(3:5))
+allocate(character(len=8) :: astr2)
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=5,kind=4) :: p4str2)
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+allocate(character(len=9,kind=4) :: a4str2)
+
+
+! integer :: ic(2:5), ic2
+
+!$omp target enter data map(alloc: dt%ic)
+!$omp target map(alloc: dt%ic)
+ if (size(dt%ic) /= 4) error stop
+ if (lbound(dt%ic, 1) /= 2) error stop
+ if (ubound(dt%ic, 1) /= 5) error stop
+ dt%ic = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic)
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: dt%ic2)
+!$omp target map(alloc: dt%ic2)
+ dt%ic2 = 42
+!$omp end target
+!$omp target exit data map(from: dt%ic2)
+if (dt%ic2 /= 42) error stop
+
+
+! character(len=11) :: ccstr(3:4), ccstr2
+
+!$omp target enter data map(alloc: dt%ccstr)
+!$omp target map(alloc: dt%ccstr)
+ if (len(dt%ccstr) /= 11) error stop
+ if (size(dt%ccstr) /= 2) error stop
+ if (lbound(dt%ccstr, 1) /= 3) error stop
+ if (ubound(dt%ccstr, 1) /= 4) error stop
+ dt%ccstr = ["12345678901", "abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr)
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+!$omp target enter data map(alloc: dt%ccstr2)
+!$omp target map(alloc: dt%ccstr2)
+ if (len(dt%ccstr2) /= 11) error stop
+ dt%ccstr2 = "ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%ccstr2)
+if (len(dt%ccstr2) /= 11) error stop
+if (dt%ccstr2 /= "ABCDEFGHIJK") error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7), cc4str2
+
+! Value check fails
+!$omp target enter data map(alloc: dt%cc4str)
+!$omp target map(alloc: dt%cc4str)
+ if (len(dt%cc4str) /= 11) error stop
+ if (size(dt%cc4str) /= 5) error stop
+ if (lbound(dt%cc4str, 1) /= 3) error stop
+ if (ubound(dt%cc4str, 1) /= 7) error stop
+ dt%cc4str = [4_"12345678901", 4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+ 4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str)
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+!$omp target enter data map(alloc: dt%cc4str2)
+!$omp target map(alloc: dt%cc4str2)
+ if (len(dt%cc4str2) /= 11) error stop
+ dt%cc4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%cc4str2)
+if (len(dt%cc4str2) /= 11) error stop
+if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! integer, pointer :: pc(:), pc2
+! allocate(dt%pc(5), dt%pc2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pc)
+!$omp target map(alloc: dt%pc)
+ if (.not. associated(dt%pc)) error stop
+ if (size(dt%pc) /= 5) error stop
+ if (lbound(dt%pc, 1) /= 1) error stop
+ if (ubound(dt%pc, 1) /= 5) error stop
+ dt%pc = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc)
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: dt%pc2)
+!$omp target map(alloc: dt%pc2)
+ if (.not. associated(dt%pc2)) error stop
+ dt%pc2 = 99
+!$omp end target
+!$omp target exit data map(from: dt%pc2)
+if (dt%pc2 /= 99) error stop
+if (.not. associated(dt%pc2)) error stop
+
+
+! character(len=:), pointer :: pcstr(:), pcstr2
+! allocate(character(len=2) :: dt%pcstr(2))
+! allocate(character(len=4) :: dt%pcstr2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pcstr)
+!$omp target map(alloc: dt%pcstr)
+ if (.not. associated(dt%pcstr)) error stop
+ if (len(dt%pcstr) /= 2) error stop
+ if (size(dt%pcstr) /= 2) error stop
+ if (lbound(dt%pcstr, 1) /= 1) error stop
+ if (ubound(dt%pcstr, 1) /= 2) error stop
+ dt%pcstr = ["01", "jk"]
+!$omp end target
+!$omp target exit data map(from: dt%pcstr)
+if (.not. associated(dt%pcstr)) error stop
+if (len(dt%pcstr) /= 2) error stop
+if (size(dt%pcstr) /= 2) error stop
+if (lbound(dt%pcstr, 1) /= 1) error stop
+if (ubound(dt%pcstr, 1) /= 2) error stop
+if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(alloc: dt%pcstr2)
+!$omp target map(alloc: dt%pcstr2)
+ if (.not. associated(dt%pcstr2)) error stop
+ if (len(dt%pcstr2) /= 4) error stop
+ dt%pcstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: dt%pcstr2)
+if (.not. associated(dt%pcstr2)) error stop
+if (len(dt%pcstr2) /= 4) error stop
+if (dt%pcstr2 /= "HIJK") error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+! allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+! structure element when other mapped elements from the same structure weren't mapped together with it
+!$omp target enter data map(alloc: dt%pc4str)
+!$omp target map(alloc: dt%pc4str)
+ if (.not. associated(dt%pc4str)) error stop
+ if (len(dt%pc4str) /= 3) error stop
+ if (size(dt%pc4str) /= 2) error stop
+ if (lbound(dt%pc4str, 1) /= 2) error stop
+ if (ubound(dt%pc4str, 1) /= 3) error stop
+ dt%pc4str = [4_"456", 4_"tzu"]
+!$omp end target
+!$omp target exit data map(from: dt%pc4str)
+if (.not. associated(dt%pc4str)) error stop
+if (len(dt%pc4str) /= 3) error stop
+if (size(dt%pc4str) /= 2) error stop
+if (lbound(dt%pc4str, 1) /= 2) error stop
+if (ubound(dt%pc4str, 1) /= 3) error stop
+if (dt%pc4str(2) /= 4_"456") error stop
+if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(alloc: dt%pc4str2)
+!$omp target map(alloc: dt%pc4str2)
+ if (.not. associated(dt%pc4str2)) error stop
+ if (len(dt%pc4str2) /= 5) error stop
+ dt%pc4str2 = 4_"98765"
+!$omp end target
+!$omp target exit data map(from: dt%pc4str2)
+if (.not. associated(dt%pc4str2)) error stop
+if (len(dt%pc4str2) /= 5) error stop
+if (dt%pc4str2 /= 4_"98765") error stop
+
+
+! integer :: ii(5), ii2
+
+!$omp target enter data map(alloc: ii)
+!$omp target map(alloc: ii)
+ if (size(ii) /= 5) error stop
+ if (lbound(ii, 1) /= 1) error stop
+ if (ubound(ii, 1) /= 5) error stop
+ ii = [-1, -2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii)
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+!$omp target enter data map(alloc: ii2)
+!$omp target map(alloc: ii2)
+ ii2 = -410
+!$omp end target
+!$omp target exit data map(from: ii2)
+if (ii2 /= -410) error stop
+
+
+! character(len=11) :: clstr(-1:1), clstr2
+
+!$omp target enter data map(alloc: clstr)
+!$omp target map(alloc: clstr)
+ if (len(clstr) /= 11) error stop
+ if (size(clstr) /= 3) error stop
+ if (lbound(clstr, 1) /= -1) error stop
+ if (ubound(clstr, 1) /= 1) error stop
+ clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr)
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+!$omp target enter data map(alloc: clstr2)
+!$omp target map(alloc: clstr2)
+ if (len(clstr2) /= 11) error stop
+ clstr2 = "ABCDEFghijk"
+!$omp end target
+!$omp target exit data map(from: clstr2)
+if (len(clstr2) /= 11) error stop
+if (clstr2 /= "ABCDEFghijk") error stop
+
+
+! character(len=11,kind=4) :: cl4str(0:3), cl4str2
+
+!$omp target enter data map(alloc: cl4str)
+!$omp target map(alloc: cl4str)
+ if (len(cl4str) /= 11) error stop
+ if (size(cl4str) /= 4) error stop
+ if (lbound(cl4str, 1) /= 0) error stop
+ if (ubound(cl4str, 1) /= 3) error stop
+ cl4str = [4_"12345678901", 4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str)
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+!$omp target enter data map(alloc: cl4str2)
+!$omp target map(alloc: cl4str2)
+ if (len(cl4str2) /= 11) error stop
+ cl4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: cl4str2)
+if (len(cl4str2) /= 11) error stop
+if (cl4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(alloc: ip)
+!$omp target map(alloc: ip)
+ if (.not. associated(ip)) error stop
+ if (size(ip) /= 5) error stop
+ if (lbound(ip, 1) /= 1) error stop
+ if (ubound(ip, 1) /= 5) error stop
+ ip = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip)
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(alloc: ip2)
+!$omp target map(alloc: ip2)
+ if (.not. associated(ip2)) error stop
+ ip2 = 99
+!$omp end target
+!$omp target exit data map(from: ip2)
+if (ip2 /= 99) error stop
+if (.not. associated(ip2)) error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(alloc: ia)
+!$omp target map(alloc: ia)
+ if (.not. allocated(ia)) error stop
+ if (size(ia) /= 8) error stop
+ if (lbound(ia, 1) /= 1) error stop
+ if (ubound(ia, 1) /= 8) error stop
+ ia = [1,2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia)
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+!$omp target enter data map(alloc: ia2)
+!$omp target map(alloc: ia2)
+ if (.not. allocated(ia2)) error stop
+ ia2 = 102
+!$omp end target
+!$omp target exit data map(from: ia2)
+if (ia2 /= 102) error stop
+if (.not. allocated(ia2)) error stop
+
+
+! character(len=:), pointer :: pstr(:), pstr2
+! allocate(character(len=2) :: pstr(-2:0))
+! allocate(character(len=4) :: pstr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: pstr)
+!$omp target map(alloc: pstr)
+ if (.not. associated(pstr)) error stop
+ if (len(pstr) /= 2) error stop
+ if (size(pstr) /= 3) error stop
+ if (lbound(pstr, 1) /= -2) error stop
+ if (ubound(pstr, 1) /= 0) error stop
+ pstr = ["01", "jk", "aq"]
+!$omp end target
+!$omp target exit data map(from: pstr)
+if (.not. associated(pstr)) error stop
+if (len(pstr) /= 2) error stop
+if (size(pstr) /= 3) error stop
+if (lbound(pstr, 1) /= -2) error stop
+if (ubound(pstr, 1) /= 0) error stop
+if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+!$omp target enter data map(alloc: pstr2)
+!$omp target map(alloc: pstr2)
+ if (.not. associated(pstr2)) error stop
+ if (len(pstr2) /= 4) error stop
+ pstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: pstr2)
+if (.not. associated(pstr2)) error stop
+if (len(pstr2) /= 4) error stop
+if (pstr2 /= "HIJK") error stop
+
+
+! character(len=:), allocatable :: astr(:), astr2
+! allocate(character(len=6) :: astr(3:5))
+! allocate(character(len=8) :: astr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: astr)
+!$omp target map(alloc: astr)
+ if (.not. allocated(astr)) error stop
+ if (len(astr) /= 6) error stop
+ if (size(astr) /= 3) error stop
+ if (lbound(astr, 1) /= 3) error stop
+ if (ubound(astr, 1) /= 5) error stop
+ astr = ["01db45", "jk$D%S", "zutg47"]
+!$omp end target
+!$omp target exit data map(from: astr)
+if (.not. allocated(astr)) error stop
+if (len(astr) /= 6) error stop
+if (size(astr) /= 3) error stop
+if (lbound(astr, 1) /= 3) error stop
+if (ubound(astr, 1) /= 5) error stop
+if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(alloc: astr2)
+!$omp target map(alloc: astr2)
+ if (.not. allocated(astr2)) error stop
+ if (len(astr2) /= 8) error stop
+ astr2 = "HIJKhijk"
+!$omp end target
+!$omp target exit data map(from: astr2)
+if (.not. allocated(astr2)) error stop
+if (len(astr2) /= 8) error stop
+if (astr2 /= "HIJKhijk") error stop
+
+
+! character(len=:,kind=4), pointer :: p4str(:), p4str2
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+! allocate(character(len=5,kind=4) :: p4str2)
+
+! FAILS with value check
+
+!$omp target enter data map(alloc: p4str)
+!$omp target map(alloc: p4str)
+ if (.not. associated(p4str)) error stop
+ if (len(p4str) /= 3) error stop
+ if (size(p4str) /= 3) error stop
+ if (lbound(p4str, 1) /= 2) error stop
+ if (ubound(p4str, 1) /= 4) error stop
+ p4str(:) = [4_"f85", 4_"8af", 4_"A%F"]
+!$omp end target
+!$omp target exit data map(from: p4str)
+if (.not. associated(p4str)) error stop
+if (len(p4str) /= 3) error stop
+if (size(p4str) /= 3) error stop
+if (lbound(p4str, 1) /= 2) error stop
+if (ubound(p4str, 1) /= 4) error stop
+if (p4str(2) /= 4_"f85") error stop
+if (p4str(3) /= 4_"8af") error stop
+if (p4str(4) /= 4_"A%F") error stop
+
+!$omp target enter data map(alloc: p4str2)
+!$omp target map(alloc: p4str2)
+ if (.not. associated(p4str2)) error stop
+ if (len(p4str2) /= 5) error stop
+ p4str2 = 4_"9875a"
+!$omp end target
+!$omp target exit data map(from: p4str2)
+if (.not. associated(p4str2)) error stop
+if (len(p4str2) /= 5) error stop
+if (p4str2 /= 4_"9875a") error stop
+
+
+! character(len=:,kind=4), allocatable :: a4str(:), a4str2
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+! allocate(character(len=9,kind=4) :: a4str2)
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+!$omp target enter data map(alloc: a4str)
+!$omp target map(alloc: a4str)
+ if (.not. allocated(a4str)) error stop
+ if (len(a4str) /= 7) error stop
+ if (size(a4str) /= 6) error stop
+ if (lbound(a4str, 1) /= -2) error stop
+ if (ubound(a4str, 1) /= 3) error stop
+ ! See PR fortran/107508 why '(:)' is required
+ a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!$omp end target
+!$omp target exit data map(from: a4str)
+if (.not. allocated(a4str)) error stop
+if (len(a4str) /= 7) error stop
+if (size(a4str) /= 6) error stop
+if (lbound(a4str, 1) /= -2) error stop
+if (ubound(a4str, 1) /= 3) error stop
+if (a4str(-2) /= 4_"sf456aq") error stop
+if (a4str(-1) /= 4_"3dtzu24") error stop
+if (a4str(0) /= 4_"_4fh7sm") error stop
+if (a4str(1) /= 4_"=ff85s7") error stop
+if (a4str(2) /= 4_"j=8af4d") error stop
+if (a4str(3) /= 4_".,A%Fsz") error stop
+
+!$omp target enter data map(alloc: a4str2)
+!$omp target map(alloc: a4str2)
+ if (.not. allocated(a4str2)) error stop
+ if (len(a4str2) /= 9) error stop
+ a4str2 = 4_"98765a23d"
+!$omp end target
+!$omp target exit data map(from: a4str2)
+if (.not. allocated(a4str2)) error stop
+if (len(a4str2) /= 9) error stop
+if (a4str2 /= 4_"98765a23d") error stop
+
+
+deallocate(dt%pc, dt%pc2)
+deallocate(dt%pcstr)
+deallocate(dt%pcstr2)
+
+deallocate(dt%pc4str)
+deallocate(dt%pc4str2)
+
+deallocate(ip, ip2, ia, ia2)
+deallocate(pstr)
+deallocate(pstr2)
+deallocate(astr)
+deallocate(astr2)
+
+deallocate(p4str)
+deallocate(p4str2)
+deallocate(a4str)
+deallocate(a4str2)
+end
--- /dev/null
+! Check that 'map((to)from:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+ integer :: ic(2:5), ic2
+ character(len=11) :: ccstr(3:4), ccstr2
+ character(len=11,kind=4) :: cc4str(3:7), cc4str2
+ integer, pointer :: pc(:), pc2
+ character(len=:), pointer :: pcstr(:), pcstr2
+ character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+end type t
+
+type(t) :: dt
+
+integer :: ii(5), ii2
+character(len=11) :: clstr(-1:1), clstr2
+character(len=11,kind=4) :: cl4str(0:3), cl4str2
+integer, pointer :: ip(:), ip2
+integer, allocatable :: ia(:), ia2
+character(len=:), pointer :: pstr(:), pstr2
+character(len=:), allocatable :: astr(:), astr2
+character(len=:,kind=4), pointer :: p4str(:), p4str2
+character(len=:,kind=4), allocatable :: a4str(:), a4str2
+
+allocate(dt%pc(5), dt%pc2)
+allocate(character(len=2) :: dt%pcstr(2))
+allocate(character(len=4) :: dt%pcstr2)
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+allocate(ip(5), ip2, ia(8), ia2)
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=4) :: pstr2)
+allocate(character(len=6) :: astr(3:5))
+allocate(character(len=8) :: astr2)
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=5,kind=4) :: p4str2)
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+allocate(character(len=9,kind=4) :: a4str2)
+
+
+! integer :: ic(2:5), ic2
+
+!$omp target enter data map(tofrom: dt%ic)
+!$omp target map(from: dt%ic)
+ if (size(dt%ic) /= 4) error stop
+ if (lbound(dt%ic, 1) /= 2) error stop
+ if (ubound(dt%ic, 1) /= 5) error stop
+ dt%ic = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic)
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: dt%ic2)
+!$omp target map(from: dt%ic2)
+ dt%ic2 = 42
+!$omp end target
+!$omp target exit data map(from: dt%ic2)
+if (dt%ic2 /= 42) error stop
+
+
+! character(len=11) :: ccstr(3:4), ccstr2
+
+!$omp target enter data map(tofrom: dt%ccstr)
+!$omp target map(from: dt%ccstr)
+ if (len(dt%ccstr) /= 11) error stop
+ if (size(dt%ccstr) /= 2) error stop
+ if (lbound(dt%ccstr, 1) /= 3) error stop
+ if (ubound(dt%ccstr, 1) /= 4) error stop
+ dt%ccstr = ["12345678901", "abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr)
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+!$omp target enter data map(tofrom: dt%ccstr2)
+!$omp target map(from: dt%ccstr2)
+ if (len(dt%ccstr2) /= 11) error stop
+ dt%ccstr2 = "ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%ccstr2)
+if (len(dt%ccstr2) /= 11) error stop
+if (dt%ccstr2 /= "ABCDEFGHIJK") error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7), cc4str2
+
+! Value check fails
+!$omp target enter data map(tofrom: dt%cc4str)
+!$omp target map(from: dt%cc4str)
+ if (len(dt%cc4str) /= 11) error stop
+ if (size(dt%cc4str) /= 5) error stop
+ if (lbound(dt%cc4str, 1) /= 3) error stop
+ if (ubound(dt%cc4str, 1) /= 7) error stop
+ dt%cc4str = [4_"12345678901", 4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+ 4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str)
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+!$omp target enter data map(tofrom: dt%cc4str2)
+!$omp target map(from: dt%cc4str2)
+ if (len(dt%cc4str2) /= 11) error stop
+ dt%cc4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: dt%cc4str2)
+if (len(dt%cc4str2) /= 11) error stop
+if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! integer, pointer :: pc(:), pc2
+! allocate(dt%pc(5), dt%pc2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pc)
+!$omp target map(from: dt%pc)
+ if (.not. associated(dt%pc)) error stop
+ if (size(dt%pc) /= 5) error stop
+ if (lbound(dt%pc, 1) /= 1) error stop
+ if (ubound(dt%pc, 1) /= 5) error stop
+ dt%pc = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc)
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: dt%pc2)
+!$omp target map(from: dt%pc2)
+ if (.not. associated(dt%pc2)) error stop
+ dt%pc2 = 99
+!$omp end target
+!$omp target exit data map(from: dt%pc2)
+if (dt%pc2 /= 99) error stop
+if (.not. associated(dt%pc2)) error stop
+
+
+! character(len=:), pointer :: pcstr(:), pcstr2
+! allocate(character(len=2) :: dt%pcstr(2))
+! allocate(character(len=4) :: dt%pcstr2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pcstr)
+!$omp target map(from: dt%pcstr)
+ if (.not. associated(dt%pcstr)) error stop
+ if (len(dt%pcstr) /= 2) error stop
+ if (size(dt%pcstr) /= 2) error stop
+ if (lbound(dt%pcstr, 1) /= 1) error stop
+ if (ubound(dt%pcstr, 1) /= 2) error stop
+ dt%pcstr = ["01", "jk"]
+!$omp end target
+!$omp target exit data map(from: dt%pcstr)
+if (.not. associated(dt%pcstr)) error stop
+if (len(dt%pcstr) /= 2) error stop
+if (size(dt%pcstr) /= 2) error stop
+if (lbound(dt%pcstr, 1) /= 1) error stop
+if (ubound(dt%pcstr, 1) /= 2) error stop
+if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(tofrom: dt%pcstr2)
+!$omp target map(from: dt%pcstr2)
+ if (.not. associated(dt%pcstr2)) error stop
+ if (len(dt%pcstr2) /= 4) error stop
+ dt%pcstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: dt%pcstr2)
+if (.not. associated(dt%pcstr2)) error stop
+if (len(dt%pcstr2) /= 4) error stop
+if (dt%pcstr2 /= "HIJK") error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:), pc4str2
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+! allocate(character(len=5,kind=4) :: dt%pc4str2)
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(tofrom: dt%pc4str)
+!$omp target map(from: dt%pc4str)
+ if (.not. associated(dt%pc4str)) error stop
+ if (len(dt%pc4str) /= 3) error stop
+ if (size(dt%pc4str) /= 2) error stop
+ if (lbound(dt%pc4str, 1) /= 2) error stop
+ if (ubound(dt%pc4str, 1) /= 3) error stop
+ dt%pc4str = [4_"456", 4_"tzu"]
+!$omp end target
+!$omp target exit data map(from: dt%pc4str)
+if (.not. associated(dt%pc4str)) error stop
+if (len(dt%pc4str) /= 3) error stop
+if (size(dt%pc4str) /= 2) error stop
+if (lbound(dt%pc4str, 1) /= 2) error stop
+if (ubound(dt%pc4str, 1) /= 3) error stop
+if (dt%pc4str(2) /= 4_"456") error stop
+if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+!$omp target enter data map(tofrom: dt%pc4str2)
+!$omp target map(from: dt%pc4str2)
+ if (.not. associated(dt%pc4str2)) error stop
+ if (len(dt%pc4str2) /= 5) error stop
+ dt%pc4str2 = 4_"98765"
+!$omp end target
+!$omp target exit data map(from: dt%pc4str2)
+if (.not. associated(dt%pc4str2)) error stop
+if (len(dt%pc4str2) /= 5) error stop
+if (dt%pc4str2 /= 4_"98765") error stop
+
+
+! integer :: ii(5), ii2
+
+!$omp target enter data map(tofrom: ii)
+!$omp target map(from: ii)
+ if (size(ii) /= 5) error stop
+ if (lbound(ii, 1) /= 1) error stop
+ if (ubound(ii, 1) /= 5) error stop
+ ii = [-1, -2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii)
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+!$omp target enter data map(tofrom: ii2)
+!$omp target map(from: ii2)
+ ii2 = -410
+!$omp end target
+!$omp target exit data map(from: ii2)
+if (ii2 /= -410) error stop
+
+
+! character(len=11) :: clstr(-1:1), clstr2
+
+!$omp target enter data map(tofrom: clstr)
+!$omp target map(from: clstr)
+ if (len(clstr) /= 11) error stop
+ if (size(clstr) /= 3) error stop
+ if (lbound(clstr, 1) /= -1) error stop
+ if (ubound(clstr, 1) /= 1) error stop
+ clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr)
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+!$omp target enter data map(tofrom: clstr2)
+!$omp target map(from: clstr2)
+ if (len(clstr2) /= 11) error stop
+ clstr2 = "ABCDEFghijk"
+!$omp end target
+!$omp target exit data map(from: clstr2)
+if (len(clstr2) /= 11) error stop
+if (clstr2 /= "ABCDEFghijk") error stop
+
+
+! character(len=11,kind=4) :: cl4str(0:3), cl4str2
+
+!$omp target enter data map(tofrom: cl4str)
+!$omp target map(from: cl4str)
+ if (len(cl4str) /= 11) error stop
+ if (size(cl4str) /= 4) error stop
+ if (lbound(cl4str, 1) /= 0) error stop
+ if (ubound(cl4str, 1) /= 3) error stop
+ cl4str = [4_"12345678901", 4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str)
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+!$omp target enter data map(tofrom: cl4str2)
+!$omp target map(from: cl4str2)
+ if (len(cl4str2) /= 11) error stop
+ cl4str2 = 4_"ABCDEFGHIJK"
+!$omp end target
+!$omp target exit data map(from: cl4str2)
+if (len(cl4str2) /= 11) error stop
+if (cl4str2 /= 4_"ABCDEFGHIJK") error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(tofrom: ip)
+!$omp target map(from: ip)
+ if (.not. associated(ip)) error stop
+ if (size(ip) /= 5) error stop
+ if (lbound(ip, 1) /= 1) error stop
+ if (ubound(ip, 1) /= 5) error stop
+ ip = [11, 22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip)
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+!$omp target enter data map(tofrom: ip2)
+!$omp target map(from: ip2)
+ if (.not. associated(ip2)) error stop
+ ip2 = 99
+!$omp end target
+!$omp target exit data map(from: ip2)
+if (ip2 /= 99) error stop
+if (.not. associated(ip2)) error stop
+
+
+! allocate(ip(5), ip2, ia(8), ia2)
+
+!$omp target enter data map(tofrom: ia)
+!$omp target map(from: ia)
+ if (.not. allocated(ia)) error stop
+ if (size(ia) /= 8) error stop
+ if (lbound(ia, 1) /= 1) error stop
+ if (ubound(ia, 1) /= 8) error stop
+ ia = [1,2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia)
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+!$omp target enter data map(tofrom: ia2)
+!$omp target map(from: ia2)
+ if (.not. allocated(ia2)) error stop
+ ia2 = 102
+!$omp end target
+!$omp target exit data map(from: ia2)
+if (ia2 /= 102) error stop
+if (.not. allocated(ia2)) error stop
+
+
+! character(len=:), pointer :: pstr(:), pstr2
+! allocate(character(len=2) :: pstr(-2:0))
+! allocate(character(len=4) :: pstr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: pstr)
+!$omp target map(from: pstr)
+ if (.not. associated(pstr)) error stop
+ if (len(pstr) /= 2) error stop
+ if (size(pstr) /= 3) error stop
+ if (lbound(pstr, 1) /= -2) error stop
+ if (ubound(pstr, 1) /= 0) error stop
+ pstr = ["01", "jk", "aq"]
+!$omp end target
+!$omp target exit data map(from: pstr)
+if (.not. associated(pstr)) error stop
+if (len(pstr) /= 2) error stop
+if (size(pstr) /= 3) error stop
+if (lbound(pstr, 1) /= -2) error stop
+if (ubound(pstr, 1) /= 0) error stop
+if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+!$omp target enter data map(tofrom: pstr2)
+!$omp target map(from: pstr2)
+ if (.not. associated(pstr2)) error stop
+ if (len(pstr2) /= 4) error stop
+ pstr2 = "HIJK"
+!$omp end target
+!$omp target exit data map(from: pstr2)
+if (.not. associated(pstr2)) error stop
+if (len(pstr2) /= 4) error stop
+if (pstr2 /= "HIJK") error stop
+
+
+! character(len=:), allocatable :: astr(:), astr2
+! allocate(character(len=6) :: astr(3:5))
+! allocate(character(len=8) :: astr2)
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: astr)
+!$omp target map(from: astr)
+ if (.not. allocated(astr)) error stop
+ if (len(astr) /= 6) error stop
+ if (size(astr) /= 3) error stop
+ if (lbound(astr, 1) /= 3) error stop
+ if (ubound(astr, 1) /= 5) error stop
+ astr = ["01db45", "jk$D%S", "zutg47"]
+!$omp end target
+!$omp target exit data map(from: astr)
+if (.not. allocated(astr)) error stop
+if (len(astr) /= 6) error stop
+if (size(astr) /= 3) error stop
+if (lbound(astr, 1) /= 3) error stop
+if (ubound(astr, 1) /= 5) error stop
+if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+
+! libgomp: nvptx_alloc error: out of memory
+
+!$omp target enter data map(tofrom: astr2)
+!$omp target map(from: astr2)
+ if (.not. allocated(astr2)) error stop
+ if (len(astr2) /= 8) error stop
+ astr2 = "HIJKhijk"
+!$omp end target
+!$omp target exit data map(from: astr2)
+if (.not. allocated(astr2)) error stop
+if (len(astr2) /= 8) error stop
+if (astr2 /= "HIJKhijk") error stop
+
+
+! character(len=:,kind=4), pointer :: p4str(:), p4str2
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+! allocate(character(len=5,kind=4) :: p4str2)
+
+! FAILS with value check
+
+!$omp target enter data map(tofrom: p4str)
+!$omp target map(from: p4str)
+ if (.not. associated(p4str)) error stop
+ if (len(p4str) /= 3) error stop
+ if (size(p4str) /= 3) error stop
+ if (lbound(p4str, 1) /= 2) error stop
+ if (ubound(p4str, 1) /= 4) error stop
+ p4str(:) = [4_"f85", 4_"8af", 4_"A%F"]
+!$omp end target
+!$omp target exit data map(from: p4str)
+if (.not. associated(p4str)) error stop
+if (len(p4str) /= 3) error stop
+if (size(p4str) /= 3) error stop
+if (lbound(p4str, 1) /= 2) error stop
+if (ubound(p4str, 1) /= 4) error stop
+if (p4str(2) /= 4_"f85") error stop
+if (p4str(3) /= 4_"8af") error stop
+if (p4str(4) /= 4_"A%F") error stop
+
+!$omp target enter data map(tofrom: p4str2)
+!$omp target map(from: p4str2)
+ if (.not. associated(p4str2)) error stop
+ if (len(p4str2) /= 5) error stop
+ p4str2 = 4_"9875a"
+!$omp end target
+!$omp target exit data map(from: p4str2)
+if (.not. associated(p4str2)) error stop
+if (len(p4str2) /= 5) error stop
+if (p4str2 /= 4_"9875a") error stop
+
+
+! character(len=:,kind=4), allocatable :: a4str(:), a4str2
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+! allocate(character(len=9,kind=4) :: a4str2)
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+!$omp target enter data map(tofrom: a4str)
+!$omp target map(from: a4str)
+ if (.not. allocated(a4str)) error stop
+ if (len(a4str) /= 7) error stop
+ if (size(a4str) /= 6) error stop
+ if (lbound(a4str, 1) /= -2) error stop
+ if (ubound(a4str, 1) /= 3) error stop
+ ! See PR fortran/107508 why '(:)' is required
+ a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!$omp end target
+!$omp target exit data map(from: a4str)
+if (.not. allocated(a4str)) error stop
+if (len(a4str) /= 7) error stop
+if (size(a4str) /= 6) error stop
+if (lbound(a4str, 1) /= -2) error stop
+if (ubound(a4str, 1) /= 3) error stop
+if (a4str(-2) /= 4_"sf456aq") error stop
+if (a4str(-1) /= 4_"3dtzu24") error stop
+if (a4str(0) /= 4_"_4fh7sm") error stop
+if (a4str(1) /= 4_"=ff85s7") error stop
+if (a4str(2) /= 4_"j=8af4d") error stop
+if (a4str(3) /= 4_".,A%Fsz") error stop
+
+!$omp target enter data map(tofrom: a4str2)
+!$omp target map(from: a4str2)
+ if (.not. allocated(a4str2)) error stop
+ if (len(a4str2) /= 9) error stop
+ a4str2 = 4_"98765a23d"
+!$omp end target
+!$omp target exit data map(from: a4str2)
+if (.not. allocated(a4str2)) error stop
+if (len(a4str2) /= 9) error stop
+if (a4str2 /= 4_"98765a23d") error stop
+
+
+deallocate(dt%pc, dt%pc2)
+deallocate(dt%pcstr)
+deallocate(dt%pcstr2)
+
+deallocate(dt%pc4str)
+deallocate(dt%pc4str2)
+
+deallocate(ip, ip2, ia, ia2)
+deallocate(pstr)
+deallocate(pstr2)
+deallocate(astr)
+deallocate(astr2)
+
+deallocate(p4str)
+deallocate(p4str2)
+deallocate(a4str)
+deallocate(a4str2)
+end
--- /dev/null
+! Check that 'map(alloc:' properly works with
+! - deferred-length character strings
+! - arrays with array descriptors
+! For those, the array descriptor / string length must be mapped with 'to:'
+
+program main
+implicit none
+
+type t
+ integer :: ic(2:5)
+ character(len=11) :: ccstr(3:4)
+ character(len=11,kind=4) :: cc4str(3:7)
+ integer, pointer :: pc(:)
+ character(len=:), pointer :: pcstr(:)
+ character(len=:,kind=4), pointer :: pc4str(:)
+end type t
+
+type(t) :: dt
+
+integer :: ii(5)
+character(len=11) :: clstr(-1:1)
+character(len=11,kind=4) :: cl4str(0:3)
+integer, pointer :: ip(:)
+integer, allocatable :: ia(:)
+character(len=:), pointer :: pstr(:)
+character(len=:), allocatable :: astr(:)
+character(len=:,kind=4), pointer :: p4str(:)
+character(len=:,kind=4), allocatable :: a4str(:)
+
+allocate(dt%pc(5))
+allocate(character(len=2) :: dt%pcstr(2))
+
+allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+
+allocate(ip(5), ia(8))
+allocate(character(len=2) :: pstr(-2:0))
+allocate(character(len=6) :: astr(3:5))
+
+allocate(character(len=3,kind=4) :: p4str(2:4))
+allocate(character(len=7,kind=4) :: a4str(-2:3))
+
+
+! integer :: ic(2:5)
+
+!$omp target enter data map(alloc: dt%ic(3:5))
+dt%ic(2) = 22
+!$omp target map(alloc: dt%ic(3:5))
+ if (size(dt%ic) /= 4) error stop
+ if (lbound(dt%ic, 1) /= 2) error stop
+ if (ubound(dt%ic, 1) /= 5) error stop
+ dt%ic(3:5) = [33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%ic(3:5))
+if (size(dt%ic) /= 4) error stop
+if (lbound(dt%ic, 1) /= 2) error stop
+if (ubound(dt%ic, 1) /= 5) error stop
+if (any (dt%ic /= [22, 33, 44, 55])) error stop
+
+! character(len=11) :: ccstr(3:4)
+
+!$omp target enter data map(alloc: dt%ccstr(4:4))
+dt%ccstr(3) = "12345678901"
+!$omp target map(alloc: dt%ccstr(4:4))
+ if (len(dt%ccstr) /= 11) error stop
+ if (size(dt%ccstr) /= 2) error stop
+ if (lbound(dt%ccstr, 1) /= 3) error stop
+ if (ubound(dt%ccstr, 1) /= 4) error stop
+ dt%ccstr(4:4) = ["abcdefghijk"]
+!$omp end target
+!$omp target exit data map(from: dt%ccstr(4:4))
+if (len(dt%ccstr) /= 11) error stop
+if (size(dt%ccstr) /= 2) error stop
+if (lbound(dt%ccstr, 1) /= 3) error stop
+if (ubound(dt%ccstr, 1) /= 4) error stop
+if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop
+
+
+! character(len=11,kind=4) :: cc4str(3:7)
+
+! Value check fails
+!$omp target enter data map(alloc: dt%cc4str(4:7))
+dt%cc4str(3) = 4_"12345678901"
+!$omp target map(alloc: dt%cc4str(4:7))
+ if (len(dt%cc4str) /= 11) error stop
+ if (size(dt%cc4str) /= 5) error stop
+ if (lbound(dt%cc4str, 1) /= 3) error stop
+ if (ubound(dt%cc4str, 1) /= 7) error stop
+ dt%cc4str(4:7) = [4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4.", &
+ 4_"45ngwj56sj2"]
+!$omp end target
+!$omp target exit data map(from: dt%cc4str(4:7))
+if (len(dt%cc4str) /= 11) error stop
+if (size(dt%cc4str) /= 5) error stop
+if (lbound(dt%cc4str, 1) /= 3) error stop
+if (ubound(dt%cc4str, 1) /= 7) error stop
+if (dt%cc4str(3) /= 4_"12345678901") error stop
+if (dt%cc4str(4) /= 4_"abcdefghijk") error stop
+if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop
+if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop
+if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop
+
+! integer, pointer :: pc(:)
+! allocate(dt%pc(5))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+!$omp target enter data map(alloc: dt%pc(2:5))
+dt%pc(1) = 11
+!$omp target map(alloc: dt%pc(2:5))
+ if (.not. associated(dt%pc)) error stop
+ if (size(dt%pc) /= 5) error stop
+ if (lbound(dt%pc, 1) /= 1) error stop
+ if (ubound(dt%pc, 1) /= 5) error stop
+ dt%pc(2:5) = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: dt%pc(2:5))
+if (.not. associated(dt%pc)) error stop
+if (size(dt%pc) /= 5) error stop
+if (lbound(dt%pc, 1) /= 1) error stop
+if (ubound(dt%pc, 1) /= 5) error stop
+if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop
+
+! character(len=:), pointer :: pcstr(:)
+! allocate(character(len=2) :: dt%pcstr(2))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+
+! FIXME: Disabled befause of PR108837
+!
+!!$omp target enter data map(alloc: dt%pcstr(2:2))
+!dt%pcstr(1) = "01"
+!!$omp target map(alloc: dt%pcstr(2:2))
+! if (.not. associated(dt%pcstr)) error stop
+! if (len(dt%pcstr) /= 2) error stop
+! if (size(dt%pcstr) /= 2) error stop
+! if (lbound(dt%pcstr, 1) /= 1) error stop
+! if (ubound(dt%pcstr, 1) /= 2) error stop
+! dt%pcstr(2:2) = ["jk"]
+!!$omp end target
+!!$omp target exit data map(from: dt%pcstr(2:2))
+!if (.not. associated(dt%pcstr)) error stop
+!if (len(dt%pcstr) /= 2) error stop
+!if (size(dt%pcstr) /= 2) error stop
+!if (lbound(dt%pcstr, 1) /= 1) error stop
+!if (ubound(dt%pcstr, 1) /= 2) error stop
+!if (any (dt%pcstr /= ["01", "jk"])) error stop
+
+
+! character(len=:,kind=4), pointer :: pc4str(:)
+! allocate(character(len=3,kind=4) :: dt%pc4str(2:3))
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00
+! structure element when other mapped elements from the same structure weren't mapped together with it
+
+! FIXME: Disabled befause of PR108837
+!
+!!$omp target enter data map(alloc: dt%pc4str(3:3))
+!dt%pc4str(2) = 4_"456"
+!!$omp target map(alloc: dt%pc4str(3:3))
+! if (.not. associated(dt%pc4str)) error stop
+! if (len(dt%pc4str) /= 3) error stop
+! if (size(dt%pc4str) /= 2) error stop
+! if (lbound(dt%pc4str, 1) /= 2) error stop
+! if (ubound(dt%pc4str, 1) /= 3) error stop
+! dt%pc4str(3:3) = [4_"tzu"]
+!!$omp end target
+!!$omp target exit data map(from: dt%pc4str(3:3))
+!if (.not. associated(dt%pc4str)) error stop
+!if (len(dt%pc4str) /= 3) error stop
+!if (size(dt%pc4str) /= 2) error stop
+!if (lbound(dt%pc4str, 1) /= 2) error stop
+!if (ubound(dt%pc4str, 1) /= 3) error stop
+!if (dt%pc4str(2) /= 4_"456") error stop
+!if (dt%pc4str(3) /= 4_"tzu") error stop
+
+! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01
+
+! integer :: ii(5)
+
+!$omp target enter data map(alloc: ii(2:5))
+ii(1) = -1
+!$omp target map(alloc: ii(2:5))
+ if (size(ii) /= 5) error stop
+ if (lbound(ii, 1) /= 1) error stop
+ if (ubound(ii, 1) /= 5) error stop
+ ii(2:5) = [-2, -3, -4, -5]
+!$omp end target
+!$omp target exit data map(from: ii(2:5))
+if (size(ii) /= 5) error stop
+if (lbound(ii, 1) /= 1) error stop
+if (ubound(ii, 1) /= 5) error stop
+if (any (ii /= [-1, -2, -3, -4, -5])) error stop
+
+
+! character(len=11) :: clstr(-1:1)
+
+!$omp target enter data map(alloc: clstr(0:1))
+clstr(-1) = "12345678901"
+!$omp target map(alloc: clstr(0:1))
+ if (len(clstr) /= 11) error stop
+ if (size(clstr) /= 3) error stop
+ if (lbound(clstr, 1) /= -1) error stop
+ if (ubound(clstr, 1) /= 1) error stop
+ clstr(0:1) = ["abcdefghijk", "ABCDEFGHIJK"]
+!$omp end target
+!$omp target exit data map(from: clstr(0:1))
+if (len(clstr) /= 11) error stop
+if (size(clstr) /= 3) error stop
+if (lbound(clstr, 1) /= -1) error stop
+if (ubound(clstr, 1) /= 1) error stop
+if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop
+
+! character(len=11,kind=4) :: cl4str(0:3)
+
+!$omp target enter data map(alloc: cl4str(1:3))
+cl4str(0) = 4_"12345678901"
+!$omp target map(alloc: cl4str(1:3))
+ if (len(cl4str) /= 11) error stop
+ if (size(cl4str) /= 4) error stop
+ if (lbound(cl4str, 1) /= 0) error stop
+ if (ubound(cl4str, 1) /= 3) error stop
+ cl4str(1:3) = [4_"abcdefghijk", &
+ 4_"qerftcea6ds", 4_"a1f9g37ga4."]
+!$omp end target
+!$omp target exit data map(from: cl4str(1:3))
+if (len(cl4str) /= 11) error stop
+if (size(cl4str) /= 4) error stop
+if (lbound(cl4str, 1) /= 0) error stop
+if (ubound(cl4str, 1) /= 3) error stop
+if (cl4str(0) /= 4_"12345678901") error stop
+if (cl4str(1) /= 4_"abcdefghijk") error stop
+if (cl4str(2) /= 4_"qerftcea6ds") error stop
+if (cl4str(3) /= 4_"a1f9g37ga4.") error stop
+
+
+! allocate(ip(5), ia(8))
+
+!$omp target enter data map(alloc: ip(2:5))
+ip(1) = 11
+!$omp target map(alloc: ip(2:5))
+ if (.not. associated(ip)) error stop
+ if (size(ip) /= 5) error stop
+ if (lbound(ip, 1) /= 1) error stop
+ if (ubound(ip, 1) /= 5) error stop
+ ip(2:5) = [22, 33, 44, 55]
+!$omp end target
+!$omp target exit data map(from: ip(2:5))
+if (.not. associated(ip)) error stop
+if (size(ip) /= 5) error stop
+if (lbound(ip, 1) /= 1) error stop
+if (ubound(ip, 1) /= 5) error stop
+if (any (ip /= [11, 22, 33, 44, 55])) error stop
+
+! allocate(ip(5), ia(8))
+
+!$omp target enter data map(alloc: ia(2:8))
+ia(1) = 1
+!$omp target map(alloc: ia(2:8))
+ if (.not. allocated(ia)) error stop
+ if (size(ia) /= 8) error stop
+ if (lbound(ia, 1) /= 1) error stop
+ if (ubound(ia, 1) /= 8) error stop
+ ia(2:8) = [2,3,4,5,6,7,8]
+!$omp end target
+!$omp target exit data map(from: ia(2:8))
+if (.not. allocated(ia)) error stop
+if (size(ia) /= 8) error stop
+if (lbound(ia, 1) /= 1) error stop
+if (ubound(ia, 1) /= 8) error stop
+if (any (ia /= [1,2,3,4,5,6,7,8])) error stop
+
+
+! character(len=:), pointer :: pstr(:)
+! allocate(character(len=2) :: pstr(-2:0))
+
+! libgomp: nvptx_alloc error: out of memory
+
+! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR
+!
+!!$omp target enter data map(alloc: pstr(-1:0))
+!pstr(-2) = "01"
+!!$omp target map(alloc: pstr(-1:0))
+! if (.not. associated(pstr)) error stop
+! if (len(pstr) /= 2) error stop
+! if (size(pstr) /= 3) error stop
+! if (lbound(pstr, 1) /= -2) error stop
+! if (ubound(pstr, 1) /= 0) error stop
+! pstr(-1:0) = ["jk", "aq"]
+!!$omp end target
+!!$omp target exit data map(from: pstr(-1:0))
+!if (.not. associated(pstr)) error stop
+!if (len(pstr) /= 2) error stop
+!if (size(pstr) /= 3) error stop
+!if (lbound(pstr, 1) /= -2) error stop
+!if (ubound(pstr, 1) /= 0) error stop
+!if (any (pstr /= ["01", "jk", "aq"])) error stop
+
+
+! character(len=:), allocatable :: astr(:)
+! allocate(character(len=6) :: astr(3:5))
+
+! libgomp: nvptx_alloc error: out of memory
+
+! FIXME
+!!$omp target enter data map(alloc: astr(4:5))
+!astr(3) = "01db45"
+!!$omp target map(alloc: astr(4:5))
+! if (.not. allocated(astr)) error stop
+! if (len(astr) /= 6) error stop
+! if (size(astr) /= 3) error stop
+! if (lbound(astr, 1) /= 3) error stop
+! if (ubound(astr, 1) /= 5) error stop
+!!! astr(4:5) = ["jk$D%S", "zutg47"]
+!!$omp end target
+!!!$omp target exit data map(from: astr(4:5))
+!!if (.not. allocated(astr)) error stop
+!!!if (len(astr) /= 6) error stop
+!if (size(astr) /= 3) error stop
+!if (lbound(astr, 1) /= 3) error stop
+!if (ubound(astr, 1) /= 5) error stop
+!if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop
+!
+
+! character(len=:,kind=4), pointer :: p4str(:)
+! allocate(character(len=3,kind=4) :: p4str(2:4))
+
+! FAILS with value check
+
+! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR
+!
+!!$omp target enter data map(alloc: p4str(3:4))
+!p4str(2) = 4_"f85"
+!!$omp target map(alloc: p4str(3:4))
+! if (.not. associated(p4str)) error stop
+! if (len(p4str) /= 3) error stop
+! if (size(p4str) /= 3) error stop
+! if (lbound(p4str, 1) /= 2) error stop
+! if (ubound(p4str, 1) /= 4) error stop
+! p4str(3:4) = [4_"8af", 4_"A%F"]
+!!$omp end target
+!!$omp target exit data map(from: p4str(3:4))
+!if (.not. associated(p4str)) error stop
+!if (len(p4str) /= 3) error stop
+!if (size(p4str) /= 3) error stop
+!if (lbound(p4str, 1) /= 2) error stop
+!if (ubound(p4str, 1) /= 4) error stop
+!if (p4str(2) /= 4_"f85") error stop
+!if (p4str(3) /= 4_"8af") error stop
+!if (p4str(4) /= 4_"A%F") error stop
+
+! character(len=:,kind=4), allocatable :: a4str(:)
+! allocate(character(len=7,kind=4) :: a4str(-2:3))
+
+! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped
+
+! FIXME: Disabled befause of PR108838
+!!$omp target enter data map(alloc: a4str(-1:3))
+!!a4str(-2) = 4_"sf456aq"
+!!$omp target map(alloc: a4str(-1:3))
+! if (.not. allocated(a4str)) error stop
+! if (len(a4str) /= 7) error stop
+! if (size(a4str) /= 6) error stop
+! if (lbound(a4str, 1) /= -2) error stop
+! if (ubound(a4str, 1) /= 3) error stop
+! a4str(-1:3) = [4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"]
+!!$omp end target
+!!$omp target exit data map(from: a4str(-1:3))
+!if (.not. allocated(a4str)) error stop
+!if (len(a4str) /= 7) error stop
+!if (size(a4str) /= 6) error stop
+!if (lbound(a4str, 1) /= -2) error stop
+!if (ubound(a4str, 1) /= 3) error stop
+!if (a4str(-2) /= 4_"sf456aq") error stop
+!if (a4str(-1) /= 4_"3dtzu24") error stop
+!if (a4str(0) /= 4_"_4fh7sm") error stop
+!if (a4str(1) /= 4_"=ff85s7") error stop
+!if (a4str(2) /= 4_"j=8af4d") error stop
+!if (a4str(3) /= 4_".,A%Fsz") error stop
+
+deallocate(dt%pc)
+deallocate(dt%pcstr)
+
+deallocate(dt%pc4str)
+
+deallocate(ip, ia)
+deallocate(pstr)
+deallocate(astr)
+
+deallocate(p4str)
+deallocate(a4str)
+end
--- /dev/null
+module m
+ implicit none
+ character(len=:), allocatable :: strA(:), strA2
+ character(len=:), pointer :: strP(:), strP2
+ !$omp declare target enter(strA,strA2,strP,strP2)
+contains
+ subroutine opt_map(str1, str2, str3)
+ character(len=:), allocatable :: str1, str2, str3, str4
+ optional :: str2, str3
+
+ if (.not.present(str2)) error stop
+ if (present(str3)) error stop
+
+ !$omp target map(str1,str2,str3,str4)
+ if (allocated(str1)) error stop
+ if (allocated(str2)) error stop
+ if (present(str3)) error stop
+ if (allocated(str4)) error stop
+ !$omp end target
+ end
+ subroutine call_opt()
+ character(len=:), allocatable :: str1, str2
+ call opt_map(str1, str2)
+ end
+ subroutine test
+ !$omp declare target
+ if (.not. allocated(strA)) error stop
+ !if (.not. allocated(strA2)) error stop
+ if (.not. associated(strP)) error stop
+ !if (.not. associated(strP2)) error stop
+
+ ! ensure length was updated as well
+ if (len(strA) /= 3) error stop
+ if (len(strA2) /= 5) error stop
+ if (len(strP) /= 4) error stop
+ if (len(strP2) /= 8) error stop
+! if (any (strA /= ['Hav', 'e f', 'un!'])) error stop
+! if (strA2 /= 'Hello') error stop
+! if (any (strP /= ['abcd', 'efgh', 'ijkl'])) error stop
+! if (strP2 /= 'TestCase') error stop
+!
+! strA = ['123', '456', '789']
+! strA2 = 'World'
+! strP = ['ABCD', 'EFGH', 'IJKL']
+! strP2 = 'Passed!!'
+ end
+end
+
+program main
+ use m
+ implicit none
+ call call_opt
+
+ strA = ['Hav', 'e f', 'un!']
+ strA2 = 'Hello'
+ allocate(character(len=4) :: strP(3))
+ strP = ['abcd', 'efgh', 'ijkl']
+ allocate(character(len=8) :: strP2)
+ strP2 = 'TestCase'
+
+ !$omp target enter data map(always, to: strA, strA2)
+ !$omp target enter data map(to: strP, strP2)
+ !$omp target
+ call test()
+ !$omp end target
+ !$omp target exit data map(always, from: strA, strA2, strP, strP2)
+
+ if (len(strA) /= 3) error stop
+ if (len(strA2) /= 5) error stop
+ if (len(strP) /= 4) error stop
+ if (len(strP2) /= 8) error stop
+! if (any (strA /= ['123', '456', '789'])) error stop
+! if (strA2 /= 'World') error stop
+! if (any(strP /= ['ABCD', 'EFGH', 'IJKL'])) error stop
+! if (strP2 /= 'Passed!!') error stop
+
+! deallocate(strP, strP2, strA, strA2)
+end
projects / thirdparty / gcc.git / commitdiff
