Message ID | c5b43e02-d1d5-e7cf-c11c-6daf1e8f33c5@codesourcery.com |
---|---|
State | New |
Headers | show |
Series | [OpenMP,Fortran] Add structure/derived-type element mapping | expand |
*PING* On 6/24/20 7:32 PM, Tobias Burnus 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. > > 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) > > Comments, remarks, suggestions? > Otherwise: OK for the trunk? > > Tobias > ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
On Wed, Jun 24, 2020 at 07:32:09PM +0200, Tobias Burnus wrote: > Comments, remarks, suggestions? > Otherwise: OK for the trunk? LGTM, thanks. > [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. Jakub
Hi Julian, Tobias! On 2020-06-24T19:32:09+0200, Tobias Burnus <tobias@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
On 6/24/20 7:32 PM, Tobias Burnus wrote: > While testing, I encountered two bugs, one relating to kind=4 > character string (patch pending review; PR95837) > not part of testcase) As that PR has been committed, I updated the testcase to check character(kind=4) as well. (I also removed the unused variables to silence -Wall warnings when running manually.) Committed as obvious. Tobias ----------------- Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander Walter
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 <tobias@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
[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