Re: [PATCH 1/2] OpenMP/Fortran: Map intermediate array descriptor [PR120505]
Paul-Antoine Arras <[email protected]>
| Newsgroups | gmane.comp.gcc.patches,gmane.comp.gcc.fortran |
|---|---|
| Message-ID | <[email protected]> |
Hi Tobias,
Thanks for your review. Here is attached a revised version of the patch.
On 11/02/2026 12:24, Tobias Burnus wrote:
[...]
> * * *
>
> Some dump tests still assume that pointers (or rather: array bounds) are
> 64 bit wide (descriptor size). I think it should handle 32bit as well.
> ('-m32')
>
> This seems to be fixed in gcc/testsuite/gfortran.dg/gomp/map-
> subarray-5.f90,
> but not in gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90.
>
> Did you run with RUNTESTFLAGS="--target_board=unix'{-m64,-m32}'" to check?
>
> I think the following needs to be also adapted ('pointer set, len: 64'):
>
> gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90:
>
> +! { dg-final { scan-tree-dump-times { #pragma omp target enter data
> map\(alloc:\*\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data
> \[len: 0\] \[runtime_implicit\]\) map\(to:chunk\.tiles \[pointer set,
> len: 64\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\)
> chunk\.tiles\.data \[bias: [0-9]+\]\) } 1 "original" } }
> ...
> +! { dg-final { scan-tree-dump-times { #pragma omp target exit data
> map\(release:chunk\.tiles \[pointer set, len: 64\]\) map\(attach_detach:
> \(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: [0-9]+
> \]\) } 1 "original" } }
>
> I think it is 36 instead of 64 with -m32 (not rechecked with this patch).
It should be (?:36|64) indeed. Fixed in map-subarray-3.f90 as well now.
> * * *
>
> Regarding:
>
> /* Sanity check: the standalone attach node will not work if we
> have
> an "enter data" operation (because for those, variables need
> to be
> mapped separately and attach nodes must be grouped together
> with the
>
>> @@ -13416,7 +13449,9 @@ omp_build_struct_sibling_lists (enum tree_code
>> code,
>> base they attach to). We should only have created the
>> ATTACH_DETACH node after GOMP_MAP_STRUCT for a target region, so
>> this should never be true. */
>> - gcc_assert ((region_type & ORT_TARGET) != 0);
>> + // This is no longer true. See zlas in gomp_map_vars_internal
>> + // (libgomp/target.c).
>> + // gcc_assert ((region_type & ORT_TARGET) != 0);
>
> The background of the comment is that for 'target data' and 'target',
> all mapped variables
> are mapped and unmapped in one step – and can be combined. For 'target
> enter/exit data',
> that's not the case while for 'target enter data map(to: a,b)' both 'a'
> and 'b' are mapped,
> either of them could be separately from the other unmapped. Thus, for
> 'target'/'target data'
> all map variables are visible to gomp_map_vars_internal – but for
> 'target enter/exit data',
> there are multiple separate calls to gomp_map_vars_internal and the
> splitting might not always
> be done in an optimal way.
>
> ZLAs are zero-length arrays - and 'a[:0]' and 'a[:5]' are handled
> differently, one as
> ZLA pointer attachment and one was map followed by an attach. If the
> size is not known at
> compile time ('a[:N]'), a nop map (nullptr, zero-size) plus an
> attachment can be generated.
> The code mentioned in libgomp is:
>
>> case GOMP_MAP_ATTACH:
>> case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
>> ...
>> bool zlas
>> = ((kind & typemask)
>> ==
>> GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
>> /* For 'target enter data', the map clauses are
>> split;
>> however, for more complex code with struct and
>> pointer members, the mapping and the attach
>> can end up
>> in different sets; or the wrong mapping with
>> the
>> attach. As there is no way to know whether a
>> size
>> zero like 'var->ptr[i][:0]' happend in the
>> same
>> directive or not, the not-attached check is now
>> fully silenced for 'enter data'. */
>> if (openmp_p && (pragma_kind &
>> GOMP_MAP_VARS_ENTER_DATA))
>> zlas = true;
>> if (!gomp_attach_pointer (devicep, aq, mem_map, n,
>> (uintptr_t)
>> hostaddrs[i], sizes[i],
>> cbufp, zlas, !openmp_p))
>> {
>> /* Pointee not found; that's an error
>> except for
>> map(var[:n]) with n == 0; the compiler
>> adds a
>> runtime condition such that for those
>> the kind is
>> always GOMP_MAP_ZERO_LEN_ARRAY_SECTION. */
> and the latter has:
>> if (!allow_zero_length_array_sections && fail_if_not_found)
>> {
>> gomp_mutex_unlock (&devicep->lock);
>> gomp_fatal ("pointer target not mapped for attach");
>
> * * *
>
> Back to the assert:
>
>> ATTACH_DETACH node after GOMP_MAP_STRUCT for a target region, so
>> this should never be true. */
>> - gcc_assert ((region_type & ORT_TARGET) != 0);
>> + // This is no longer true. See zlas in gomp_map_vars_internal
>> + // (libgomp/target.c).
>> + // gcc_assert ((region_type & ORT_TARGET) != 0);
>
> The commented 'gcc_assert' implies that this patch now generates such
> attach also for target (enter) data :-)
>
> I wonder whether we can do any better than just removing the comment and
> the check for good. - Or, alternatively, leave a stub explanation that
> there
> are issues with doing it here - even thought it kind of works.
What about setting a bool when an attach node needs adjustment and then
checking it in the assert (and updating the comment accordingly ofc)?
See the amended patch in attachment.
> * * *
>
> Back to the latest patch:
>
>> if (wholestruct)
>> {
>> + tree desc = OMP_CLAUSE_CHAIN (*(*wholestruct)->grp_start);
>> + if (desc != NULL_TREE && omp_map_clause_descriptor_p (desc))
>> + goto next;
>> *mapped_by_group = *wholestruct;
>> return true;
>> }
>
> I know that we tend to underdocument things, but I think adding
> a short comment would help in the future.
I know it all too well :/
Added comment thusly.
> * * *
>
>> + /* Find each attach node whose bias needs to be adjusted and move
>> it to the
>> + * group containing its pointee, right after the struct node. */
>> + FOR_EACH_VEC_ELT (*groups, i, grp)
>> ...
>> + && OMP_CLAUSE_MAP_KIND (grp->grp_end) == GOMP_MAP_ATTACH_DETACH
>> + && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end))
>
> I think the comment needs to be expanded why the bias is adjusted by
> moving it
> right after the struct node. The added code doesn't do it (as one would
> expect
> naively from reading the code) as it happens indirectly.
Expanded the comment as suggested.
> * * *
>
>> +/* Helper function for gfc_trans_omp_clauses. */
>> +
>> +static bool
>> +gfc_map_array_descriptor (
>> + tree &node, tree &node2, tree &node3, tree &node4, tree descr, bool
>> openacc,
>> + location_t map_loc, stmtblock_t *block, gfc_exec_op op,
>> gfc_omp_namelist *n,
>> + hash_map<gfc_symbol *, gfc_omp_namelist *> *&sym_rooted_nl, gfc_se se,
>> + gfc_omp_clauses *clauses, bool mid_desc_p)
>
> Can you at least roughly describe (in a comment) what that function
> actually does? Including stating what the return value means.
Added description in comment.
> * * *
>
>> (1) map (alloc: *(struct tile_type[0:] * restrict) chunk.tiles.data
>> [len: 0])
>> (2) map (to: chunk.tiles [pointer set, len: 64])
>> (3) map (attach_detach: (struct tile_type[0:] * restrict)
>> chunk.tiles.data
>> [bias: -1])
>>
>> (1) will turn into a no-op at runtime because the inner component is
>> explicitly
>> to-mapped but alloc is required at compile time for attaching. (2)
>> ensures that
>> the array descriptor will be available at runtime to compute offsets
>> and strides
>> in various dimensions. The gimplifier will turn (3) into a regular
>> attach of the
>> data pointer and compute the bias.
>
> It seems as if we should remove (1) after/during gimplification if we
> know that it
> is only used intermittently. Maybe add another flag for this?
As suggested off list, created a flag OMP_CLAUSE_MAP_GIMPLE_ONLY for
that purpose.
I also added the new flags to the tree pretty printer and updated the
scan dumps.
Thanks,
--
PA
0001-OpenMP-Fortran-Map-intermediate-array-descriptor-PR1.patch
(text/x-patch, 32 KB)
From 5d4dbb700852ff220dff358267aff5526d20361e Mon Sep 17 00:00:00 2001 From: Paul-Antoine Arras <[email protected]> Date: Tue, 9 Dec 2025 19:26:39 +0100 Subject: [PATCH] OpenMP/Fortran: Map intermediate array descriptor [PR120505] Consider the following OMP directive, assuming tiles is allocatable: !$omp target enter data & !$omp map(to: chunk%tiles(1)%field%density0) & !$omp map(to: chunk%left_rcv_buffer) libgomp reports an illegal memory access error at runtime. This is because density0 is referenced through tiles, which requires its descriptor to be mapped along its content. This patch ensures that such intervening allocatable in a reference chain is properly mapped. For the above example, the frontend has to create the following three additional map clauses: (1) map (alloc: *(struct tile_type[0:] * restrict) chunk.tiles.data [len: 0]) (2) map (to: chunk.tiles [pointer set, len: 64]) (3) map (attach_detach: (struct tile_type[0:] * restrict) chunk.tiles.data [bias: -1]) (1) is required by the gimplifier for attaching but will be removed at the end of the pass; the inner component is explicitly to-mapped elsewhere. (2) ensures that the array descriptor will be available at runtime to compute offsets and strides in various dimensions. The gimplifier will turn (3) into a regular attach of the data pointer and compute the bias. PR fortran/120505 gcc/fortran/ChangeLog: * trans-openmp.cc (gfc_map_array_descriptor): New function. (gfc_trans_omp_clauses): Emit map clauses for an intermediate array descriptor. gcc/ChangeLog: * gimplify.cc (omp_mapped_by_containing_struct): Handle Fortran array descriptors. (omp_build_struct_sibling_lists): Allow attach_detach bias to be adjusted on non-target regions. (gimplify_adjust_omp_clauses): Remove GIMPLE-only nodes. * tree-pretty-print.cc (dump_omp_clause): Handle OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT and OMP_CLAUSE_MAP_GIMPLE_ONLY. * tree.h (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT, OMP_CLAUSE_MAP_GIMPLE_ONLY): Define. libgomp/ChangeLog: * testsuite/libgomp.fortran/map-subarray-11.f90: New test. * testsuite/libgomp.fortran/map-subarray-13.f90: New test. * testsuite/libgomp.fortran/map-alloc-present-2.f90: New file. gcc/testsuite/ChangeLog: * gfortran.dg/gomp/map-subarray-3.f90: New test. * gfortran.dg/gomp/map-subarray-5.f90: New test. --- gcc/fortran/trans-openmp.cc | 362 ++++++++++-------- gcc/gimplify.cc | 54 ++- .../gfortran.dg/gomp/map-subarray-3.f90 | 49 +++ .../gfortran.dg/gomp/map-subarray-5.f90 | 53 +++ gcc/tree-pretty-print.cc | 13 +- gcc/tree.h | 9 + .../libgomp.fortran/map-alloc-present-2.f90 | 50 +++ .../libgomp.fortran/map-subarray-11.f90 | 56 +++ .../libgomp.fortran/map-subarray-13.f90 | 50 +++ 9 files changed, 540 insertions(+), 156 deletions(-) create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/map-alloc-present-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-11.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-13.f90 diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 86f83e6bb1f..27004a02edb 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3531,6 +3531,179 @@ get_symbol_rooted_namelist (hash_map<gfc_symbol *, return NULL; } +/* Helper function for gfc_trans_omp_clauses. Adjust existing and create new + * map nodes for derived-type component array descriptors. Return true if the + * mapping has to be dropped. */ + +static bool +gfc_map_array_descriptor ( + tree &node, tree &node2, tree &node3, tree &node4, tree descr, bool openacc, + location_t map_loc, stmtblock_t *block, gfc_exec_op op, gfc_omp_namelist *n, + hash_map<gfc_symbol *, gfc_omp_namelist *> *&sym_rooted_nl, gfc_se se, + gfc_omp_clauses *clauses, bool mid_desc_p) +{ + tree type = TREE_TYPE (descr); + tree ptr = gfc_conv_descriptor_data_get (descr); + ptr = build_fold_indirect_ref (ptr); + OMP_CLAUSE_DECL (node) = ptr; + int rank = GFC_TYPE_ARRAY_RANK (type); + OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, descr, rank); + tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + + gomp_map_kind map_kind = OMP_CLAUSE_MAP_KIND (node); + if (GOMP_MAP_COPY_TO_P (map_kind) || map_kind == GOMP_MAP_ALLOC) + { + if (mid_desc_p) + { + /* For an intermediate descriptor, the pointee (i.e. the actual array + * content) is mapped in a separate set of nodes. This ALLOC is only + * emitted to comply with the group layout expected by the gimplifier. + */ + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); + OMP_CLAUSE_SIZE (node) = size_int (0); + OMP_CLAUSE_MAP_GIMPLE_ONLY (node) = 1; + } + else + map_kind + = ((GOMP_MAP_ALWAYS_P (map_kind) || gfc_expr_attr (n->expr).pointer) + ? GOMP_MAP_ALWAYS_TO + : GOMP_MAP_TO); + } + else if (n->u.map.op == OMP_MAP_RELEASE || n->u.map.op == OMP_MAP_DELETE) + ; + else if (op == EXEC_OMP_TARGET_EXIT_DATA || op == EXEC_OACC_EXIT_DATA) + map_kind = GOMP_MAP_RELEASE; + else if (mid_desc_p) + { + /* For an intermediate descriptor, the pointee (i.e. the actual array + * content) is mapped in a separate set of nodes. This ALLOC is only + * emitted to comply with the group layout expected by the gimplifier. */ + OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC); + OMP_CLAUSE_SIZE (node) = size_int (0); + OMP_CLAUSE_MAP_GIMPLE_ONLY (node) = 1; + } + 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 (map_loc, 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); + + node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP); + if (map_kind == GOMP_MAP_RELEASE || map_kind == GOMP_MAP_DELETE) + { + OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); + OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1; + } + else + OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET); + OMP_CLAUSE_DECL (node2) = descr; + OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); + + if (!openacc) + { + if (n->expr->ts.type == BT_DERIVED + && n->expr->ts.u.derived->attr.alloc_comp) + { + /* Save array descriptor for use + in gfc_omp_deep_mapping{,_p,_cnt}; force + evaluate to ensure that it is + not gimplified + is a decl. */ + tree tmp = OMP_CLAUSE_SIZE (node); + tree var = gfc_create_var (TREE_TYPE (tmp), NULL); + gfc_add_modify_loc (map_loc, block, var, tmp); + OMP_CLAUSE_SIZE (node) = var; + gfc_allocate_lang_decl (var); + GFC_DECL_SAVED_DESCRIPTOR (var) = descr; + } + + gfc_omp_namelist *n2 = clauses->lists[OMP_LIST_MAP]; + + /* If we don't have a mapping of a smaller part + of the array -- or we can't prove that we do + statically -- set this flag. If there is a + mapping of a smaller part of the array after + all, this will turn into a no-op at + runtime. */ + OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1; + + bool sym_based; + n2 = get_symbol_rooted_namelist (sym_rooted_nl, n, n2, &sym_based); + + bool drop_mapping = false; + + for (; n2 != NULL; n2 = n2->next) + { + if ((!sym_based && n == n2) || (sym_based && n == n2->u2.duplicate_of) + || !n2->expr) + continue; + + if (!gfc_omp_expr_prefix_same (n->expr, n2->expr)) + continue; + + gfc_ref *ref1 = n->expr->ref; + gfc_ref *ref2 = n2->expr->ref; + + /* We know ref1 and ref2 overlap. We're + interested in whether ref2 describes a + smaller part of the array than ref1, which + we already know refers to the full + array. */ + + while (ref1->next && ref2->next) + { + ref1 = ref1->next; + ref2 = ref2->next; + } + + if (ref2->next + || (ref2->type == REF_ARRAY + && (ref2->u.ar.type == AR_ELEMENT + || (ref2->u.ar.type == AR_SECTION)))) + { + drop_mapping = true; + break; + } + } + if (drop_mapping) + return true; + } + + if (mid_desc_p && GOMP_MAP_COPY_FROM_P (OMP_CLAUSE_MAP_KIND (node))) + node = NULL_TREE; + + node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH_DETACH); + OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (descr); + /* 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); + if (mid_desc_p) + OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (node3) = 1; + + return false; +} + static tree gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, locus where, bool declare_simd = false, @@ -3544,6 +3717,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, enum omp_clause_code clause_code; gfc_omp_namelist *prev = NULL; gfc_se se; + vec<gfc_symbol *> descriptors = vNULL; if (clauses == NULL) return NULL_TREE; @@ -4645,6 +4819,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { gfc_init_se (&se, NULL); se.expr = gfc_maybe_dereference_var (n->sym, decl); + tree mid_descr = NULL_TREE; + gfc_ref *midref = NULL; for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next) { @@ -4654,6 +4830,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, conv_parent_component_references (&se, ref); gfc_conv_component_ref (&se, ref); + if (!mid_descr + && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (se.expr))) + { + mid_descr = se.expr; + midref = ref; + } } else if (ref->type == REF_ARRAY) { @@ -4807,156 +4989,11 @@ 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 type = TREE_TYPE (inner); - tree ptr = gfc_conv_descriptor_data_get (inner); - ptr = build_fold_indirect_ref (ptr); - OMP_CLAUSE_DECL (node) = ptr; - 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)); - map_kind = OMP_CLAUSE_MAP_KIND (node); - if (GOMP_MAP_COPY_TO_P (map_kind) - || map_kind == GOMP_MAP_ALLOC) - map_kind = ((GOMP_MAP_ALWAYS_P (map_kind) - || gfc_expr_attr (n->expr).pointer) - ? GOMP_MAP_ALWAYS_TO : GOMP_MAP_TO); - else if (n->u.map.op == OMP_MAP_RELEASE - || n->u.map.op == OMP_MAP_DELETE) - ; - else if (op == EXEC_OMP_TARGET_EXIT_DATA - || op == EXEC_OACC_EXIT_DATA) - map_kind = GOMP_MAP_RELEASE; - 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 (map_loc, 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); - node2 = build_omp_clause (map_loc, OMP_CLAUSE_MAP); - if (map_kind == GOMP_MAP_RELEASE - || map_kind == GOMP_MAP_DELETE) - { - OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); - OMP_CLAUSE_RELEASE_DESCRIPTOR (node2) = 1; - } - else - OMP_CLAUSE_SET_MAP_KIND (node2, - GOMP_MAP_TO_PSET); - OMP_CLAUSE_DECL (node2) = inner; - OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - if (!openacc) - { - if (n->expr->ts.type == BT_DERIVED - && n->expr->ts.u.derived->attr.alloc_comp) - { - /* Save array descriptor for use - in gfc_omp_deep_mapping{,_p,_cnt}; force - evaluate to ensure that it is - not gimplified + is a decl. */ - tree tmp = OMP_CLAUSE_SIZE (node); - tree var = gfc_create_var (TREE_TYPE (tmp), - NULL); - gfc_add_modify_loc (map_loc, block, - var, tmp); - OMP_CLAUSE_SIZE (node) = var; - gfc_allocate_lang_decl (var); - GFC_DECL_SAVED_DESCRIPTOR (var) = inner; - } - - gfc_omp_namelist *n2 - = clauses->lists[OMP_LIST_MAP]; - - /* If we don't have a mapping of a smaller part - of the array -- or we can't prove that we do - statically -- set this flag. If there is a - mapping of a smaller part of the array after - all, this will turn into a no-op at - runtime. */ - OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (node) = 1; - - bool sym_based; - n2 = get_symbol_rooted_namelist (sym_rooted_nl, - n, n2, - &sym_based); - - bool drop_mapping = false; - - for (; n2 != NULL; n2 = n2->next) - { - if ((!sym_based && n == n2) - || (sym_based && n == n2->u2.duplicate_of) - || !n2->expr) - continue; - - if (!gfc_omp_expr_prefix_same (n->expr, - n2->expr)) - continue; - - gfc_ref *ref1 = n->expr->ref; - gfc_ref *ref2 = n2->expr->ref; - - /* We know ref1 and ref2 overlap. We're - interested in whether ref2 describes a - smaller part of the array than ref1, which - we already know refers to the full - array. */ - - while (ref1->next && ref2->next) - { - ref1 = ref1->next; - ref2 = ref2->next; - } - - if (ref2->next - || (ref2->type == REF_ARRAY - && (ref2->u.ar.type == AR_ELEMENT - || (ref2->u.ar.type - == AR_SECTION)))) - { - drop_mapping = true; - break; - } - } - if (drop_mapping) - continue; - } - node3 = build_omp_clause (map_loc, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, - GOMP_MAP_ATTACH_DETACH); - 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); + bool drop_mapping = gfc_map_array_descriptor ( + node, node2, node3, node4, inner, openacc, map_loc, + block, op, n, sym_rooted_nl, se, clauses, false); + if (drop_mapping) + continue; } else OMP_CLAUSE_DECL (node) = inner; @@ -4972,6 +5009,31 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } else gcc_unreachable (); + + /* Map intermediate array descriptor. */ + if (!openacc && mid_descr != NULL_TREE && mid_descr != inner + && !descriptors.contains (midref->u.c.sym)) + { + descriptors.safe_push (midref->u.c.sym); + + tree node1 = copy_node (node); + tree node2 = NULL_TREE; + tree node3 = NULL_TREE; + tree node4 = NULL_TREE; + gfc_map_array_descriptor (node1, node2, node3, node4, + mid_descr, openacc, map_loc, + block, op, n, sym_rooted_nl, se, + clauses, true); + + if (node1 != NULL_TREE) + omp_clauses = gfc_trans_add_clause (node1, omp_clauses); + if (node2 != NULL_TREE) + omp_clauses = gfc_trans_add_clause (node2, omp_clauses); + if (node3 != NULL_TREE) + omp_clauses = gfc_trans_add_clause (node3, omp_clauses); + if (node4 != NULL_TREE) + omp_clauses = gfc_trans_add_clause (node4, omp_clauses); + } } else sorry_at (gfc_get_location (&n->where), "unhandled expression"); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index a27dc46d88c..bbbe6e9f7c2 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -11511,9 +11511,16 @@ omp_mapped_by_containing_struct (hash_map<tree_operand_hash_no_se, } if (wholestruct) { + /* An intermediate descriptor should not match here because the + * pointee is actually not mapped by this group -- it is just a + * zero-length alloc. */ + tree desc = OMP_CLAUSE_CHAIN (*(*wholestruct)->grp_start); + if (desc != NULL_TREE && omp_map_clause_descriptor_p (desc)) + goto next; *mapped_by_group = *wholestruct; return true; } + next: decl = wsdecl; } @@ -13390,6 +13397,39 @@ omp_build_struct_sibling_lists (enum tree_code code, tail = added_tail; } + /* Find each attach node whose bias needs to be adjusted and move it to the + * group containing its pointee, right after the struct node, so that it can + * be picked up by the adjustment code further down in this function. */ + bool attach_bias_needs_adjustment; + attach_bias_needs_adjustment = false; + FOR_EACH_VEC_ELT (*groups, i, grp) + { + tree c = *grp->grp_start; + if (c != NULL && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT + && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_TO_PSET + && OMP_CLAUSE_MAP_KIND (grp->grp_end) == GOMP_MAP_ATTACH_DETACH + && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end)) + { + OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end) = 0; + attach_bias_needs_adjustment = true; + tree *cp; + for (cp = &OMP_CLAUSE_CHAIN (c); cp != NULL; + cp = &OMP_CLAUSE_CHAIN (*cp)) + if (*cp == grp->grp_end) + { + c = *cp; + break; + } + + tree base = OMP_CLAUSE_DECL (c); + gcc_assert (TREE_CODE (base) == NOP_EXPR); + base = build_fold_indirect_ref (base); + tree *struct_node = struct_map_to_clause->get (base); + omp_siblist_move_node_after (c, cp, &OMP_CLAUSE_CHAIN (*struct_node)); + } + } + /* Now we have finished building the struct sibling lists, reprocess newly-added "attach" nodes: we need the address of the first mapped element of each struct sibling list for the bias of the attach @@ -13414,9 +13454,11 @@ omp_build_struct_sibling_lists (enum tree_code code, an "enter data" operation (because for those, variables need to be mapped separately and attach nodes must be grouped together with the base they attach to). We should only have created the - ATTACH_DETACH node after GOMP_MAP_STRUCT for a target region, so - this should never be true. */ - gcc_assert ((region_type & ORT_TARGET) != 0); + ATTACH_DETACH node either after GOMP_MAP_STRUCT for a target region + or for an intermediate descriptor that needs adjustment -- so this + should never be true. */ + gcc_assert ((region_type & ORT_TARGET) != 0 + || attach_bias_needs_adjustment); /* This is the first sorted node in the struct sibling list. Use it to recalculate the correct bias to use. @@ -15832,6 +15874,11 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, break; case OMP_CLAUSE_MAP: + if (OMP_CLAUSE_MAP_GIMPLE_ONLY (c)) + { + remove = true; + goto end_adjust_omp_map_clause; + } decl = OMP_CLAUSE_DECL (c); if (!grp_end) { @@ -16030,6 +16077,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p, /* Fallthrough. */ case OMP_TARGET_EXIT_DATA: OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DETACH); + OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (c) = 0; break; case OACC_UPDATE: /* An "attach/detach" operation on an update directive diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90 new file mode 100644 index 00000000000..aebe713d0f5 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-3.f90 @@ -0,0 +1,49 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } + +! PR fortran/120505 + +! Check that mapping nested allocatable DT components triggers required +! additional mappings for the outer array descriptor. + +module m + + type field_type + real(kind=8), allocatable :: density0(:,:), density1(:,:) + end type field_type + + type tile_type + type(field_type) :: field + end type tile_type + + type chunk_type + real(kind=8), allocatable :: left_rcv_buffer(:) + type(tile_type), allocatable :: tiles(:) + end type chunk_type + + type(chunk_type) :: chunk + +end + +use m + +allocate(chunk%tiles(1)) +chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2]) +chunk%tiles(1)%field%density1 = reshape([5,6,7,8],[2,2]) +allocate(chunk%left_rcv_buffer(1)) + +!$omp target enter data & +!$omp map(to: chunk%tiles(1)%field%density0) & +!$omp map(to: chunk%tiles(1)%field%density1) & +!$omp map(to: chunk%left_rcv_buffer) + +! { dg-final { scan-tree-dump-times { #pragma omp target enter data map\(alloc:\*\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[len: 0\] \[runtime_implicit\] \[gimple only\]\) map\(to:chunk\.tiles \[pointer set, len: (?:36|64)\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: 0 \(needs adjustment\)\]\) } 1 "original" } } + +!$omp target exit data & +!$omp map(from: chunk%tiles(1)%field%density0) & +!$omp map(from: chunk%tiles(1)%field%density1) & +!$omp map(from: chunk%left_rcv_buffer) + +! { dg-final { scan-tree-dump-times { #pragma omp target exit data map\(release:chunk\.tiles \[pointer set, len: (?:36|64)\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: 0 \(needs adjustment\)\]\) } 1 "original" } } + +end diff --git a/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90 b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90 new file mode 100644 index 00000000000..71b6296d1fc --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-subarray-5.f90 @@ -0,0 +1,53 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +! PR fortran/120505 + +! Check that the bias into the inner derived type is correctly computed on +! target enter data. For target exit data, the bias is ignored so just check +! that detach is present. +! Pointer set lengths are checked for both 32 and 64 bits. + +module m + + type field_type + real(kind=8), allocatable :: density0(:,:), density1(:,:) + end type field_type + + type tile_type + type(field_type) :: field + end type tile_type + + type chunk_type + real(kind=8), allocatable :: left_rcv_buffer(:) + type(tile_type), allocatable :: tiles(:) + end type chunk_type + + type(chunk_type) :: chunk + +end + +use m + +allocate(chunk%tiles(1)) +chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2]) +allocate(chunk%left_rcv_buffer(1)) + +!$omp target enter data & +!$omp map(to: chunk%tiles(1)%field%density1) & +!$omp map(to: chunk%left_rcv_buffer) + +! { dg-final { scan-tree-dump-times { map\(struct_unord:MEM <struct tile_type\[0:\]> \[\(struct tile_type\[0:\] \*\)_[0-9]+\] \[len: 1\]\) map\(to:MEM <struct tile_type\[0:\]> \[\(struct tile_type\[0:\] \*\)_[0-9]+\]\[_[0-9]+\]\.field\.density1 \[pointer set, len: (?:48|88)\]\) map\(attach:chunk\.tiles\.data \[bias: _[0-9]+\]\) } 1 "gimple" } } + +!$omp target exit data & +!$omp map(from: chunk%tiles(1)%field%density1) & +!$omp map(from: chunk%left_rcv_buffer) + +! { dg-final { scan-tree-dump-times { map\(release:chunk\.tiles \[pointer set, len: (?:36|64)\]\) map\(detach:chunk\.tiles\.data \[bias: [0-9]+\]\)} 1 "gimple" } } + + +! { dg-final { scan-tree-dump-not { map\(alloc } "gimple" } } +! { dg-final { scan-tree-dump-not { gimple only } "gimple" } } +! { dg-final { scan-tree-dump-not { needs adjustment } "gimple" } } + +end diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index ca1a6f2f470..79e0a734dbb 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1200,11 +1200,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, dump_flags_t flags) } dump_generic_node (pp, OMP_CLAUSE_SIZE (clause), spc, flags, false); + if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP + && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (clause)) + pp_string (pp, " (needs adjustment)"); pp_right_bracket (pp); } - if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP - && OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause)) - pp_string (pp, " [runtime_implicit]"); + if (OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP) + { + if (OMP_CLAUSE_MAP_RUNTIME_IMPLICIT_P (clause)) + pp_string (pp, " [runtime_implicit]"); + if (OMP_CLAUSE_MAP_GIMPLE_ONLY (clause)) + pp_string (pp, " [gimple only]"); + } pp_right_paren (pp); break; diff --git a/gcc/tree.h b/gcc/tree.h index 49a443f74ed..a97688e16a9 100644 --- a/gcc/tree.h +++ b/gcc/tree.h @@ -1939,6 +1939,15 @@ class auto_suppress_location_wrappers #define OMP_CLAUSE_MAP_READONLY(NODE) \ TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) +/* Nonzero if the size (or bias) is not known by the front-end and needs to be adjusted in the middle end. */ +#define OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT(NODE) \ + TREE_CONSTANT (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + +/* Nonzero on a map clause that is only used internally by the gimplifier and + * can thus be removed at the end of the GIMPLE pass. */ +#define OMP_CLAUSE_MAP_GIMPLE_ONLY(NODE) \ + TREE_USED (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE_MAP)) + /* Same as above, for use in OpenACC cache directives. */ #define OMP_CLAUSE__CACHE__READONLY(NODE) \ TREE_READONLY (OMP_CLAUSE_SUBCODE_CHECK (NODE, OMP_CLAUSE__CACHE_)) diff --git a/libgomp/testsuite/libgomp.fortran/map-alloc-present-2.f90 b/libgomp/testsuite/libgomp.fortran/map-alloc-present-2.f90 new file mode 100644 index 00000000000..4fab8a36eb7 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-alloc-present-2.f90 @@ -0,0 +1,50 @@ +! Test the 'present' modifier with derived-type allocatable components. + +module m + implicit none + type field_type + real(kind=8), allocatable :: density0(:,:), density1(:,:) + end type field_type + + type tile_type + type(field_type) :: field + end type tile_type + + type chunk_type + real(kind=8), allocatable :: left_rcv_buffer(:) + type(tile_type), allocatable :: tiles(:) + end type chunk_type + + type(chunk_type) :: chunk +end + +use m +implicit none +allocate(chunk%tiles(1)) +chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2]) + +!$omp target enter data & +!$omp map(to: chunk%tiles(1)%field%density0) & +!$omp map(to: chunk%tiles(1)%field%density1) + +!$omp target map(present, alloc: chunk%tiles(1)%field%density0) + if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1 + if (any (chunk%tiles(1)%field%density0 /= reshape([1,2,3,4],[2,2]))) stop 1 + chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 2 +!$omp end target + +chunk%tiles(1)%field%density1 = reshape([11,22,33,44],[2,2]) + +!$omp target map(alloc: chunk%tiles(1)%field%density0) + if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1 + if (any (chunk%tiles(1)%field%density0 /= 2*reshape([1,2,3,4],[2,2]))) stop 1 + chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 7 +!$omp end target + +!$omp target exit data & +!$omp map(from: chunk%tiles(1)%field%density0) + +if (any (chunk%tiles(1)%field%density0 /= 7*2*reshape([1,2,3,4],[2,2]))) stop 1 +if (any (chunk%tiles(1)%field%density1 /= reshape([11,22,33,44],[2,2]))) stop 2 + +end diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90 new file mode 100644 index 00000000000..d81a31491bb --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-11.f90 @@ -0,0 +1,56 @@ +! { dg-do run } + +! PR fortran/120505 + +! Check that mapping nested allocatable DT components triggers required +! additional mappings for the outer array descriptor. + +module m + + type field_type + real(kind=8), allocatable :: density0(:,:), density1(:,:) + end type field_type + + type tile_type + type(field_type) :: field + end type tile_type + + type chunk_type + real(kind=8), allocatable :: left_rcv_buffer(:) + type(tile_type), allocatable :: tiles(:) + end type chunk_type + + type(chunk_type) :: chunk + +end + +use m + +allocate(chunk%tiles(1)) +chunk%tiles(1)%field%density0 = reshape([1,2,3,4],[2,2]) +chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2]) +allocate(chunk%left_rcv_buffer(1)) + +!$omp target enter data & +!$omp map(to: chunk%tiles(1)%field%density0) & +!$omp map(to: chunk%tiles(1)%field%density1) & +!$omp map(to: chunk%left_rcv_buffer) + +!$omp target + if (any (chunk%tiles(1)%field%density0 /= reshape([1,2,3,4],[2,2]))) stop 1 + if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1 + chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 + 7 + chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5 + chunk%left_rcv_buffer(1) = 42.0_8 +!$omp end target + +!$omp target exit data & +!$omp map(from: chunk%tiles(1)%field%density0) & +!$omp map(from: chunk%tiles(1)%field%density1) & +!$omp map(from: chunk%left_rcv_buffer) + +if (any (chunk%tiles(1)%field%density0 /= 7 + reshape([1,2,3,4],[2,2]))) stop 1 +if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1 +if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1 + +end diff --git a/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90 b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90 new file mode 100644 index 00000000000..13fad59415b --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/map-subarray-13.f90 @@ -0,0 +1,50 @@ +! { dg-do run } + +! PR fortran/120505 + +! Check that a nested allocatable DT component is mapped properly even when the +! first component is *not* mapped. + +module m + + type field_type + real(kind=8), allocatable :: density0(:,:), density1(:,:) + end type field_type + + type tile_type + type(field_type) :: field + end type tile_type + + type chunk_type + real(kind=8), allocatable :: left_rcv_buffer(:) + type(tile_type), allocatable :: tiles(:) + end type chunk_type + + type(chunk_type) :: chunk + +end + +use m + +allocate(chunk%tiles(1)) +chunk%tiles(1)%field%density1 = reshape([1,2,3,4],[2,2]) +allocate(chunk%left_rcv_buffer(1)) + +!$omp target enter data & +!$omp map(to: chunk%tiles(1)%field%density1) & +!$omp map(to: chunk%left_rcv_buffer) + +!$omp target + if (any (chunk%tiles(1)%field%density1 /= reshape([1,2,3,4],[2,2]))) stop 1 + chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 + 5 + chunk%left_rcv_buffer(1) = 42.0_8 +!$omp end target + +!$omp target exit data & +!$omp map(from: chunk%tiles(1)%field%density1) & +!$omp map(from: chunk%left_rcv_buffer) + +if (any (chunk%tiles(1)%field%density1 /= 5 + reshape([1,2,3,4],[2,2]))) stop 1 +if (chunk%left_rcv_buffer(1) /= 42.0_8) stop 1 + +end -- 2.51.0