Hi!
If I got my tracking right, the og10 commit
4677091db1aa9d2a52e4839812bd73f47cc5e421 "[OpenMP, Fortran] Add
structure/derived-type element mapping" regresses:
[-PASS:-]{+FAIL:+} gfortran.dg/goacc/pr70828.f90 -O
scan-tree-dump-times gimple "omp target oacc_data map\\(tofrom:MEM\\[\\(c_char
\\*\\)_[0-9]+\\] \\[len: _[0-9]+\\]\\) map\\(alloc:data \\[pointer assign, b
ias: _[0-9]+\\]\\)" 1
[-PASS:-]{+FAIL:+} gfortran.dg/goacc/pr70828.f90 -O
scan-tree-dump-times gimple "omp target oacc_parallel
map\\(force_present:MEM\\[\\(c_char \\*\\)D\\.[0-9]+\\] \\[len:
D\\.[0-9]+\\]\\) map\\(alloc:data \\[
pointer assign, bias: D\\.[0-9]+\\]\\)" 1
PASS: gfortran.dg/goacc/pr70828.f90 -O (test for excess errors)
Tobias, please have a look.
And then, the issues mentioned before (Julian, Tobias -- that's what we
talked about on the phone call earlier today):
On 2020-07-15T08:33:00+0200, I wrote:
> On 2020-06-24T19:32:09+0200, Tobias Burnus <tob...@codesourcery.com> 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
