Hi Julian, Tobias!
On 2020-06-24T19:32:09+0200, Tobias Burnus <[email protected]> wrote:
> (OpenMP 5 extends this a lot, but this is about OpenMP 4.5.
> It touches code which is also used by OpenACC's attach/detach.)
>
> @OpenACC/Julian: I think the character attach/detach for
> deferred-length strings does not work properly with OpenACC;
> I did not touch this code – but I think it needs some love.
Please file a PR.
> This code adds support for
> map(dt%comp, dt%comp2)
> where "comp" can be either a nonpointer, nonallocatable element
> scalar, array or array section. Or it can be a pointer - where
> character strings are one complication as for deferred-length
> ones, the length is stored in an extra DT component.
>
> While testing, I encountered two bugs, one relating to kind=4
> character string (patch pending review; PR95837)
> not part of testcase) and one related to deferred-length
> character strings (commented in the test case; larger issue;
> PR95868).
>
> Like always, some more tests/testcase probably would not harm.
>
> Regarding the patch:
>
> (a) openmp.c:
> This enabled component matching for 'map(' and
> piggybacks on the OpenACC code for the checks. I think that
> some additional checks might be useful – and I hope that no
> check is too strict.
> The "depend" clause was excluded as one otherwise gets a
> testsuite fails due to the is-contiguous check.
>
> (b) trans-openmp.c:
> - gfc_trans_omp_clauses now has a "bool openacc".
> - GOMP_MAP_ATTACH_DETACH is replaced by GOMP_MAP_ALWAYS_POINTER
> - For arrays, the mapping of the descriptor is squeezed before
> "node" which contains the data transfer (var.desc.data mapping
> followed by the always_pointer for the mapping).
> In this array case, the latter gets a pointless cast in order
> to prevent that for both var.desc and var.desc.data memory gets
> allocated in the struct.
> → That's also the reason the big switch table is moved up.
> - For deferred-length strings, the string-length is in an extra
> struct element (derived-type component) and will be mapped in
> addition.
> - Bugs in the previous version:
> * gfc_trans_omp_array_section for "element == true", the size
> of a pointer instead of the size of the element was mapped.
> * For string variables (with constant length) the kind=4 was
> not properly handled.
> * Allocatable scalars were not handled – missing second clause
> for the always_pointer (and attach_detach, I assume)
I understand correctly that your remark "Bugs in the previous version"
translates to "bugs still existing on releases/gcc-10 branch for OpenACC
'attach'/'detach'"? Should we thus backport to releases/gcc-10 branch
this commit 102502e32ea4e8a75d6b252ba319d09d735d9aa7 "[OpenMP, Fortran]
Add structure/derived-type element mapping", and fixup commit
524862db444b6544c6dc87c5f06f351100ecf50d "Fix goacc/finalize-1.f tree
dump-scanning for -m32"?
Grüße
Thomas
> Comments, remarks, suggestions?
> Otherwise: OK for the trunk?
> [OpenMP, Fortran] Add structure/derived-type element mapping
>
> gcc/fortran/ChangeLog:
>
> * openmp.c (gfc_match_omp_clauses): Match also derived-type
> component refs in OMP_CLAUSE_MAP.
> (resolve_omp_clauses): Resolve those.
> * trans-openmp.c (gfc_trans_omp_array_section, gfc_trans_omp_clauses):
> Handle OpenMP structure-element mapping.
> (gfc_trans_oacc_construct, gfc_trans_oacc_executable_directive,
> (gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Update
> add openacc=true in gfc_trans_omp_clauses call.
>
> gcc/testsuite/ChangeLog:
>
> * gfortran.dg/goacc/finalize-1.f: Update dump scan pattern.
> * gfortran.dg/gomp/map-1.f90: Update dg-error.
> * gfortran.dg/gomp/map-2.f90: New test.
>
>
> libgomp/ChangeLog:
>
> * testsuite/libgomp.fortran/struct-elem-map-1.f90: New test.
>
> gcc/fortran/openmp.c | 5 +-
> gcc/fortran/trans-openmp.c | 332
> +++++++++++++++------
> gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 +-
> gcc/testsuite/gfortran.dg/gomp/map-1.f90 | 35 +--
> gcc/testsuite/gfortran.dg/gomp/map-2.f90 | 6 +
> .../libgomp.fortran/struct-elem-map-1.f90 | 331 ++++++++++++++++++++
> 6 files changed, 595 insertions(+), 118 deletions(-)
>
> diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
> index e681903c7c2..7de2f6e1b1d 100644
> --- a/gcc/fortran/openmp.c
> +++ b/gcc/fortran/openmp.c
> @@ -1464,7 +1464,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const
> omp_mask mask,
> 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)
> @@ -4553,7 +4553,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses
> *omp_clauses,
>
> /* 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
> @@ -4563,6 +4563,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses
> *omp_clauses,
> 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",
> diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
> index 7e2f6256c43..3f4f06375ef 100644
> --- a/gcc/fortran/trans-openmp.c
> +++ b/gcc/fortran/trans-openmp.c
> @@ -2087,10 +2087,11 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
> 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);
>
> @@ -2099,7 +2100,8 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> gfc_omp_namelist *n,
> 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
> {
> @@ -2109,14 +2111,15 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> gfc_omp_namelist *n,
> 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)))
> @@ -2129,28 +2132,71 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> gfc_omp_namelist *n,
> 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)));
> @@ -2161,14 +2207,15 @@ gfc_trans_omp_array_section (stmtblock_t *block,
> gfc_omp_namelist *n,
> 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;
> @@ -2483,6 +2530,67 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> 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;
> @@ -2491,7 +2599,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> && 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)
> @@ -2719,8 +2827,42 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> /* 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;
> }
>
> @@ -2747,7 +2889,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> 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;
>
> @@ -2768,7 +2910,9 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> 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);
> }
> @@ -2795,32 +2939,82 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
>
> 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;
> @@ -2832,9 +3026,11 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> && 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. */
> @@ -2846,65 +3042,7 @@ gfc_trans_omp_clauses (stmtblock_t *block,
> gfc_omp_clauses *clauses,
> }
>
> 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);
> @@ -3656,7 +3794,7 @@ gfc_trans_oacc_construct (gfc_code *code)
>
> 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);
> @@ -3692,7 +3830,7 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
>
> 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);
> @@ -4517,7 +4655,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
> 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 = █
> @@ -5695,7 +5833,7 @@ gfc_trans_oacc_declare (gfc_code *code)
> 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);
> diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> index 1e2e3e94b8a..fd496968506 100644
> --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f
> @@ -20,7 +20,7 @@
> ! { 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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
>
> !$ACC EXIT DATA COPYOUT (cpo_r)
> @@ -32,6 +32,6 @@
> ! { 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: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
> END SUBROUTINE f
> diff --git a/gcc/testsuite/gfortran.dg/gomp/map-1.f90
> b/gcc/testsuite/gfortran.dg/gomp/map-1.f90
> index e78b56c8f39..831feffcc43 100644
> --- a/gcc/testsuite/gfortran.dg/gomp/map-1.f90
> +++ b/gcc/testsuite/gfortran.dg/gomp/map-1.f90
> @@ -57,18 +57,20 @@ subroutine test(aas)
> !$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
> @@ -81,29 +83,28 @@ subroutine test(aas)
>
> !$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
> diff --git a/gcc/testsuite/gfortran.dg/gomp/map-2.f90
> b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
> new file mode 100644
> index 00000000000..73c4f5a87d0
> --- /dev/null
> +++ b/gcc/testsuite/gfortran.dg/gomp/map-2.f90
> @@ -0,0 +1,6 @@
> +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
> diff --git a/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
> b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
> new file mode 100644
> index 00000000000..f18eeb90165
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90
> @@ -0,0 +1,331 @@
> +! { 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
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
