[1/2] OpenMP/Fortran: Map intermediate array descriptor [PR120505]
Commit Message
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) 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.
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.
* tree-core.h (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT): Define.
libgomp/ChangeLog:
* testsuite/libgomp.fortran/map-subarray-11.f90: New test.
* testsuite/libgomp.fortran/map-subarray-13.f90: New test.
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 | 345 ++++++++++--------
gcc/gimplify.cc | 37 +-
.../gfortran.dg/gomp/map-subarray-3.f90 | 49 +++
.../gfortran.dg/gomp/map-subarray-5.f90 | 47 +++
gcc/tree-core.h | 1 +
.../libgomp.fortran/map-subarray-11.f90 | 56 +++
.../libgomp.fortran/map-subarray-13.f90 | 50 +++
7 files changed, 434 insertions(+), 151 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-subarray-11.f90
create mode 100644 libgomp/testsuite/libgomp.fortran/map-subarray-13.f90
Comments
Hi PA,
not a review but a head up (or heads down, I am currently confused …).
Paul-Antoine Arras wrote:
> 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])
I think I need to think about this a bit more.
To have something lighter, I tried:
----------------------
integer, allocatable :: aa(:)
integer, pointer :: pp(:)
pp => null()
!$omp target enter data map(aa)
!$omp target enter data map(pp)
!$omp target map(present, alloc: aa, pp)
if (associated(pp) .or. allocated(aa)) i = 1
!$omp end target
!$omp target exit data map(pp)
!$omp target exit data map(aa)
allocate(aa, pp, source=[1,2,3])
!$omp target enter data map(pp)
!$omp target enter data map(aa)
!$omp target map(always, to: aa) map(to: pp)
! GCC + ftn: 'map(present, alloc:' -> 'aa' and 'pp' not in the present table
! if (associated(pp) .or. allocated(aa)) stop 1
! if (any (pp /= [1,2,3])) stop 1
! if (any (aa /= [1,2,3])) stop 1
pp = pp * 2
aa = aa * 3
!$omp end target
!$omp target exit data map(from: aa)
!$omp target exit data map(from: pp)
print *, aa ! 3,6,9
print *, pp ! 0 (?) w/ cray, 2,4,6 with GCC.
end
-----------------
I think the GCC result for that program (as currently written)
makes sense - but I am not sure I understand the 'present' → error not present.
I guess, I need to re-read the specification here.
If one comments the inner enter/exit for 'aa', only 'to' is
active and the result is the outer '1,2,3' - kind of makes sense
if 'aa' is not in the present table. (But should it?)
* * *
The original starting point is the following program where for
the first 'target' region, Cray ftn works (accepting the 'present')
but gfortran fails with:
The following program fails with GCC (and hopefully
correctly applied patches) with:
libgomp: Trying to map into
device [0x62d1a50..0x62d1b00) object
when [0x62d1a50..0x62d1aa8) is already mapped
which I find rather odd. If one uses 'to' instead of 'present'
(+ some ignored map type, let's pick: 'alloc'), it works with
GCC as well.
(BTW: Cray ftn rejects 'print' inside target and gives an ICE
when using 'stop'.)
* * *
Likewise for the second target region, where GCC does not
like the 'present' either. Using
'alloc: ... density0'
'always, to: density1'
it fails differently:
libgomp: cuCtxSynchronize error: an illegal memory access was encountered
However, with
'to: density0'
'always, to: density1'
the program compiles and runs past this target region.
However, at runtime, 'from:' in target exit data doesn't bring the data back
for 'density1' (but for density0) - while 'always, from' (for density1)
will cause:
libgomp: cuCtxSynchronize error: an illegal memory access was encountered
Again, this message is a bit surprising - while the failing copy back seems
to be due to 'density1' not being in the present table, I'd guess.
As Cray ftn shares the not-in-present-table behavior for the scalar
case, it is not surprising that it also uses the host value for 'density1'.
But it doesn't have the odd crash GCC has. It behaves identical for
'always, from' and 'from', contrary to GCC:
-------------------------
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(present, alloc: chunk%tiles(1)%field%density0) &
!$omp map(always, present, to: chunk%tiles(1)%field%density1)
! if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1
! if (.not. allocated(chunk%tiles(1)%field%density1)) stop 1
! if (any (chunk%tiles(1)%field%density0 /= 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 1
chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 7
chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 * 3
!$omp end target
!$omp target exit data &
!$omp map(from: chunk%tiles(1)%field%density0) &
!$omp map(from: chunk%tiles(1)%field%density1)
print *, chunk%tiles(1)%field%density0
print *, chunk%tiles(1)%field%density1
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 /= 3*reshape([11,22,33,44],[2,2]))) stop 2
end
-------------------------
* * *
Tobias,
who is now trying to understand when things are supposed to end up in the
present table and when only the data and when the pointed-to data gets
mapped.
OpenMP 6.0 added ref_ptee, ref_ptr, and ref_ptr_ptee as map modifiers, which
might help to explain some fine print a bit better as those are for mapping
the pointer target vs. the pointee. (I think some fine print might have been
fixed in TR14 or post-TR14, i.e. reading the newest version possible might
help.)
On 20/01/2026 13:17, Tobias Burnus wrote:
> The original starting point is the following program where for
> the first 'target' region, Cray ftn works (accepting the 'present')
> but gfortran fails with:
>
> The following program fails with GCC (and hopefully
> correctly applied patches) with:
> libgomp: Trying to map into
> device [0x62d1a50..0x62d1b00) object
> when [0x62d1a50..0x62d1aa8) is already mapped
>
> which I find rather odd. If one uses 'to' instead of 'present'
> (+ some ignored map type, let's pick: 'alloc'), it works with
> GCC as well.
...
> -------------------------
> 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(present, alloc: chunk%tiles(1)%field%density0) &
> !$omp map(always, present, to: chunk%tiles(1)%field%density1)
> ! if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1
> ! if (.not. allocated(chunk%tiles(1)%field%density1)) stop 1
> ! if (any (chunk%tiles(1)%field%density0 /= 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 1
> chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 7
> chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 * 3
> !$omp end target
>
> !$omp target exit data &
> !$omp map(from: chunk%tiles(1)%field%density0) &
> !$omp map(from: chunk%tiles(1)%field%density1)
>
> print *, chunk%tiles(1)%field%density0
> print *, chunk%tiles(1)%field%density1
>
> 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 /= 3*reshape([11,22,33,44],[2,2]))) stop 2
>
> end
> -------------------------
The attached diff fixes the overlapping-maps error for the testcase
above. However it then fails for `map(always, present, to:
chunk%tiles(1)%field%density1)` with:
libgomp: present clause: not present on the device (addr: 0x4bd95b0,
size: 32 (0x20), dev: 0)
This does not seem to be caused by my intermediate-descriptor patch
though, as the equivalent testcase without DT (see attachment) fails the
same way. I think I see what the problem is though and will try to come
up with a separate patch. At this point, my understanding is that the
array *descriptor* of density1 is mapped by `target enter data` but,
because it is unallocated, the array itself is not added to the present
table.
On 20/01/2026 13:17, Tobias Burnus wrote:
> * * *
>
> Likewise for the second target region, where GCC does not
> like the 'present' either. Using
>
> 'alloc: ... density0'
> 'always, to: density1'
>
> it fails differently:
> libgomp: cuCtxSynchronize error: an illegal memory access was encountered
>
> However, with
> 'to: density0'
> 'always, to: density1'
> the program compiles and runs past this target region.
>
> However, at runtime, 'from:' in target exit data doesn't bring the data back
> for 'density1' (but for density0) - while 'always, from' (for density1)
> will cause:
> libgomp: cuCtxSynchronize error: an illegal memory access was encountered
>
> Again, this message is a bit surprising - while the failing copy back seems
> to be due to 'density1' not being in the present table, I'd guess.
Here is my understanding of what GCC and libgomp currently handle this
case in turn:
1) 'target enter data'
a) maps both the descriptor and the data of density0 as it is
allocated and initialised - both refcounts are set to 1;
b) maps the descriptor of density1 but does not create storage on the
device for its data since it is still unallocated - the descriptor's
refcount is set to 1.
2) The first 'target map' runs fine because the *data* of density0 is in
the present table.
3) The second 'target map'
a) works fine for density0 as above;
b) fails for density1 because its data is not in the present table,
even though its descriptor is.
4) Assuming the present modifier is stripped and 3) does not fail,
'target exit data':
a) transfers the data of density0 back to the host;
b) upon exit from 3), the refcount of density1's descriptor is still
1 but that of its data is 0 so it gets unmapped without a chance for the
dev2host transfer to happen.
So we have two issues: the present modifier and the unmapping.
For the former, I would suggest to apply the present modifier to the
array descriptor rather than its data.
For the latter, it is not clear to me whether the OpenMP spec mandates
that both the array descriptor and the data share the same refcount or not.
> As Cray ftn shares the not-in-present-table behavior for the scalar
> case, it is not surprising that it also uses the host value for 'density1'.
> But it doesn't have the odd crash GCC has. It behaves identical for
> 'always, from' and 'from', contrary to GCC:
>
> -------------------------
> 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(present, alloc: chunk%tiles(1)%field%density0) &
> !$omp map(always, present, to: chunk%tiles(1)%field%density1)
> ! if (.not. allocated(chunk%tiles(1)%field%density0)) stop 1
> ! if (.not. allocated(chunk%tiles(1)%field%density1)) stop 1
> ! if (any (chunk%tiles(1)%field%density0 /= 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 1
> chunk%tiles(1)%field%density0 = chunk%tiles(1)%field%density0 * 7
> chunk%tiles(1)%field%density1 = chunk%tiles(1)%field%density1 * 3
> !$omp end target
>
> !$omp target exit data &
> !$omp map(from: chunk%tiles(1)%field%density0) &
> !$omp map(from: chunk%tiles(1)%field%density1)
>
> print *, chunk%tiles(1)%field%density0
> print *, chunk%tiles(1)%field%density1
>
> 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 /= 3*reshape([11,22,33,44],[2,2]))) stop 2
>
> end
> -------------------------
Hi PA,
Paul-Antoine Arras wrote:
> The attached diff fixes the overlapping-maps error for the testcase
> above. However it then fails for `map(always, present, to:
> chunk%tiles(1)%field%density1)` with:
>
> libgomp: present clause: not present on the device (addr: 0x4bd95b0,
> size: 32 (0x20), dev: 0)
With current mainline and 1/2, I see:
libgomp: Trying to map into device [0x19591a50..0x19591b00) object
when [0x19591a50..0x19591aa8) is already mapped
for the first 'target enter data':
!$omp target enter data &
!$omp map(to: chunk%tiles(1)%field%density0) &
!$omp map(to: chunk%tiles(1)%field%density1)
while with this patch, it only fails when present-always-mapping
'density1' after allocation - failing with:
libgomp: present clause: not present on the device (addr: 0x39633290,
size: 32 (0x20), dev: 0)
* * *
An obvious testcase is then the attached trimmed-down version of the test,
which does not access the unallocated density1 on the device. It seems to
me perfectly valid and sensible.
Thus – as you planned: It makes sense to fold 'present.diff' into the
1/2 patch, but not without a testcase. Voila, the attachment is one.
(I come back to the patch review of present.diff as part of the review of
1/2 itself.)
* * *
Regarding the attached testcase (present-nodt.f90), I notice that
with your follow-up patch, it behaves as follows (tested with the 1/2 patch
of this thread + current mainline): As it, it runs but the value is not copied out ("STOP 2");
using 'tofrom' instead of 'to' for 'density1', it compiles and works.
This is kind of sensible. Do you know whether there is already test coverage
for the 'tofrom' case? If not, I think a testcase should be added. [I know that
some variant it, but I lost track of the fine print and I don't feel looking
at existing testcases.]
Tobias
Hi Tobias,
Please find attached an updated version of the patch and a few comments
below.
On 06/02/2026 18:13, Tobias Burnus wrote:
> Paul-Antoine Arras wrote:
>
>> The attached diff fixes the overlapping-maps error for the testcase
>> above. However it then fails for `map(always, present, to:
>> chunk%tiles(1)%field%density1)` with:
>>
>> libgomp: present clause: not present on the device (addr: 0x4bd95b0,
>> size: 32 (0x20), dev: 0)
>
> With current mainline and 1/2, I see:
>
> libgomp: Trying to map into device [0x19591a50..0x19591b00) object
> when [0x19591a50..0x19591aa8) is already mapped
>
> for the first 'target enter data':
>
> !$omp target enter data &
> !$omp map(to: chunk%tiles(1)%field%density0) &
> !$omp map(to: chunk%tiles(1)%field%density1)
>
> while with this patch, it only fails when present-always-mapping
> 'density1' after allocation - failing with:
>
> libgomp: present clause: not present on the device (addr: 0x39633290,
> size: 32 (0x20), dev: 0)
This last error is expected and is not introduced by this patch -- it
has to do with the way the 'present' modifier is handled for
allocatables in general.
I'll try to come up with a fix in a separate patch.
For future reference, r16-7288-g1e71ff87c97fcd fixes the same issue for
bare allocatables (outside any derived type).
> * * *
>
> An obvious testcase is then the attached trimmed-down version of the test,
> which does not access the unallocated density1 on the device. It seems to
> me perfectly valid and sensible.
>
> Thus – as you planned: It makes sense to fold 'present.diff' into the
> 1/2 patch, but not without a testcase. Voila, the attachment is one.
'present.diff' is now included and 'present-dt.f90' is added as
map-alloc-present-2.f90.
> (I come back to the patch review of present.diff as part of the review of
> 1/2 itself.)
>
> * * *
>
> Regarding the attached testcase (present-nodt.f90), I notice that
> with your follow-up patch, it behaves as follows (tested with the 1/2 patch
> of this thread + current mainline): As it, it runs but the value is not
> copied out ("STOP 2");
> using 'tofrom' instead of 'to' for 'density1', it compiles and works.
As discussed elsewhere, this is due to the underspecified behaviour of
allocatable mapping (even without DT). The current libgomp
implementation treats Fortran pointers and allocatables the same, so
that pointer and pointee have different refcounts.
> This is kind of sensible. Do you know whether there is already test
> coverage
> for the 'tofrom' case? If not, I think a testcase should be added. [I
> know that
> some variant it, but I lost track of the fine print and I don't feel
> looking
> at existing testcases.]
In map-alloc-present-2.f90, I changed 'to' into 'tofrom', as the
behaviour observed with the former is controversial.
Please find attached a slightly amended version of the patch with
OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT defined as a flag rather than a
special value.
Thanks,
Hi PA,
Paul-Antoine Arras wrote:
> Please find attached a slightly amended version of the patch with
> OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT defined as a flag rather than a
> special value.
Thanks. I was about to write an email about this.
The following is copied from the draft reply to your previous email,
I think it still applies. And sorry for the slow patch review with
intervals due to getting interrupted by other things all the time.
* * *
Paul-Antoine Arras wrote:
>> With current mainline and 1/2, I see:
>>
>> libgomp: Trying to map into device [0x19591a50..0x19591b00) object
>> when [0x19591a50..0x19591aa8) is already mapped
>>
>> for the first 'target enter data':
...
> This last error is expected and is not introduced by this patch -- it
> has to do with the way the 'present' modifier is handled for
> allocatables in general.
> I'll try to come up with a fix in a separate patch.
I didn't meant to imply that it was. And thanks for looking into it.
> For future reference, r16-7288-g1e71ff87c97fcd fixes the same issue
> for bare allocatables (outside any derived type).
* * *
>> Regarding the attached testcase (present-nodt.f90), I notice that
>> with your follow-up patch, it behaves as follows (tested with the 1/2
>> patch
>> of this thread + current mainline): As it, it runs but the value is
>> not copied out ("STOP 2");
>> using 'tofrom' instead of 'to' for 'density1', it compiles and works.
>
> As discussed elsewhere, this is due to the underspecified behaviour of
> allocatable mapping (even without DT). The current libgomp
> implementation treats Fortran pointers and allocatables the same, so
> that pointer and pointee have different refcounts.
The latter seems to be an issue relative to OpenMP <= 5.2, but seems to match OpenMP 6.0+
as a result of the ref_ptr/ref_ptee changes there.
* * *
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).
* * *
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.
* * *
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.
* * *
> + /* 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.
* * *
> +/* 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.
* * *
> (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?
Tobias,
who still has to continue reviewing the rest of the patch.
Paul-Antoine Arras wrote:
> Please find attached a slightly amended version of the patch with
> OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT defined as a flag rather than a
> special value.
I misread one part of the patch – but it turned out be helpful,
nonetheless, because the testcase I created fails (see attachment).
With current mainline (w/o offloading configured), the testcase passes
(compiles and runs with host fallback).
However, with the 1/2 patch applied, it segfaults:
libgomp/testsuite/libgomp.fortran/map-subarray-13.f90:41:38:
41 | !$omp map(to: chunk%left_rcv_buffer)
| ^
internal compiler error: Segmentation fault
0x26213cd internal_error(char const*, ...)
../../../repos/gcc/gcc/diagnostic-global-context.cc:787
0x123d12f crash_signal
../../../repos/gcc/gcc/toplev.cc:325
0xec3ebe omp_build_struct_sibling_lists
../../../repos/gcc/gcc/gimplify.cc:13421
0xed14b7 gimplify_adjust_omp_clauses
../../../repos/gcc/gcc/gimplify.cc:15660
0xef11be gimplify_omp_target_update
../../../repos/gcc/gcc/gimplify.cc:19054
The testcase is a modified version of your testcase, with some additional
layers of allocatables added.
The fail is because struct_node == NULL:
13402 if (c != NULL && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
13403 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
13404 && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) == GOMP_MAP_TO_PSET
13405 && OMP_CLAUSE_MAP_KIND (grp->grp_end) == GOMP_MAP_ATTACH_DETACH
13406 && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end))
13407 {
...
13420 tree *struct_node = struct_map_to_clause->get (base);
13421 omp_siblist_move_node_after (c, cp, &OMP_CLAUSE_CHAIN (*struct_node));
(gdb) p struct_node
$11 = (tree *) 0x0
(gdb) p debug(c)
<omp_clause 0x7ffff727a2d0 constant map
op-0: <nop_expr 0x7ffff728a7a0>
op-1: <integer_cst 0x7ffff7402df8 0>
op-2:>
$7 = void
(gdb) p c == *cp
$9 = true
(gdb) p debug(base)
<indirect_ref 0x7ffff7366a40
type <array_type 0x7ffff722f0a8
type <record_type 0x7ffff7223690 tile_type2 BLK
...
arg:0 <nop_expr 0x7ffff728a7a0
arg:0 <component_ref 0x7ffff7289b70 type <pointer_type 0x7ffff7424738>
arg:0 <component_ref 0x7ffff7341210 type <record_type 0x7ffff722f000 array01_tile_type2>
arg:0 <var_decl 0x7ffff72324c0 chunk> arg:1 <field_decl 0x7ffff7230140 tiles>
* * *
Testcase attached (modified version of your testcase).
Tobias
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,
On 16/02/2026 18:49, Tobias Burnus wrote:
> Paul-Antoine Arras wrote:
>> Please find attached a slightly amended version of the patch with
>> OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT defined as a flag rather than a
>> special value.
>
> I misread one part of the patch – but it turned out be helpful,
> nonetheless, because the testcase I created fails (see attachment).
>
> With current mainline (w/o offloading configured), the testcase passes
> (compiles and runs with host fallback).
>
> However, with the 1/2 patch applied, it segfaults:
>
> libgomp/testsuite/libgomp.fortran/map-subarray-13.f90:41:38:
>
> 41 | !$omp map(to: chunk%left_rcv_buffer)
> | ^
> internal compiler error: Segmentation fault
> 0x26213cd internal_error(char const*, ...)
> ../../../repos/gcc/gcc/diagnostic-global-context.cc:787
> 0x123d12f crash_signal
> ../../../repos/gcc/gcc/toplev.cc:325
> 0xec3ebe omp_build_struct_sibling_lists
> ../../../repos/gcc/gcc/gimplify.cc:13421
> 0xed14b7 gimplify_adjust_omp_clauses
> ../../../repos/gcc/gcc/gimplify.cc:15660
> 0xef11be gimplify_omp_target_update
> ../../../repos/gcc/gcc/gimplify.cc:19054
>
>
> The testcase is a modified version of your testcase, *with some additional
> layers of allocatables added*.
Initially I crafted this patch for the "simple" case with only one
intermediate descriptor, leaving the more complex cases for later -- and
kind of forgot about it.
To support additional layers in the FE would probably mean saving in a
vector all allocatables found while walking the chain of references.
Then loop through the vector and call gfc_map_array_descriptor for each
element. Hopefully that will just work with no changes in the ME.
If so, I'll come up with a revised version of the patch. Otherwise, if
it entails much more work, what about putting sorry() there and leaving
a note in the PR?
> The fail is because struct_node == NULL:
>
> 13402 if (c != NULL && OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
> 13403 && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
> 13404 && OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) ==
> GOMP_MAP_TO_PSET
> 13405 && OMP_CLAUSE_MAP_KIND (grp->grp_end) ==
> GOMP_MAP_ATTACH_DETACH
> 13406 && OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT (grp->grp_end))
> 13407 {
> ...
> 13420 tree *struct_node = struct_map_to_clause->get (base);
> 13421 omp_siblist_move_node_after (c, cp, &OMP_CLAUSE_CHAIN
> (*struct_node));
>
>
> (gdb) p struct_node
> $11 = (tree *) 0x0
>
> (gdb) p debug(c)
> <omp_clause 0x7ffff727a2d0 constant map
> op-0: <nop_expr 0x7ffff728a7a0>
> op-1: <integer_cst 0x7ffff7402df8 0>
> op-2:>
> $7 = void
>
> (gdb) p c == *cp
> $9 = true
>
>
> (gdb) p debug(base)
> <indirect_ref 0x7ffff7366a40
> type <array_type 0x7ffff722f0a8
> type <record_type 0x7ffff7223690 tile_type2 BLK
> ...
> arg:0 <nop_expr 0x7ffff728a7a0
> arg:0 <component_ref 0x7ffff7289b70 type <pointer_type
> 0x7ffff7424738>
> arg:0 <component_ref 0x7ffff7341210 type <record_type
> 0x7ffff722f000 array01_tile_type2>
> arg:0 <var_decl 0x7ffff72324c0 chunk> arg:1 <field_decl
> 0x7ffff7230140 tiles>
>
> * * *
>
> Testcase attached (modified version of your testcase).
>
> Tobias
Hi PA,
Paul-Antoine Arras wrote:
> Initially I crafted this patch for the "simple" case with only one
> intermediate descriptor, leaving the more complex cases for later --
> and kind of forgot about it.
I have another variant – see attachment – which also gives an ICE,
but this time for:
10971 node = OMP_CLAUSE_CHAIN (node);
10972 if (node && omp_map_clause_descriptor_p (node))
10973 {
10974 gcc_assert (node != grp->grp_end);
10975 node = OMP_CLAUSE_CHAIN (node);
10976 }
as the node == grp->grp_end.
The variable is:
<omp_clause 0x7ffff7073140 map
op-0: <component_ref 0x7ffff70ddab0>
op-1: <integer_cst 0x7ffff726fd08 64>
op-2:>
<component_ref 0x7ffff70ddab0
type <record_type 0x7ffff702ae70 array01_tile_type type_1 BLK
arg:0 <var_decl 0x7ffff7032390 chunk
arg:1 <field_decl 0x7ffff702cc80 tiles type <record_type 0x7ffff702ae70 array01_tile_type>
Also this variant shouldn't ICE - and also only gives an ICE with your patch.
Without, it works with host fallback or self mapping (unsurprisingly), but fails
with actual mapping with: "is already mapped"
(I have not thought about whether that code makes sense or not,
but an ICE is definitely wrong.)
* * *
>
> To support additional layers in the FE would probably mean saving in a
> vector all allocatables found while walking the chain of references.
> Then loop through the vector and call gfc_map_array_descriptor for
> each element. Hopefully that will just work with no changes in the ME.
>
> If so, I'll come up with a revised version of the patch.
Fingers crossed!
> Otherwise, if it entails much more work, what about putting sorry()
> there and leaving a note in the PR?
Presumably, that's the best. — Although, the question is whether it
starts to make more sense to open a new PR for it. Additionally, I fear
that no one will pick up the follow up task if not done properly now.
(Whether in this patch or in a soon-to-follow follow-up patch.)
Tobias
Hi Tobias,
Here is a revised version of the patch.
On 17/02/2026 11:57, Tobias Burnus wrote:
> Hi PA,
>
> Paul-Antoine Arras wrote:
>> Initially I crafted this patch for the "simple" case with only one
>> intermediate descriptor, leaving the more complex cases for later --
>> and kind of forgot about it.
>
> I have another variant – see attachment – which also gives an ICE,
> but this time for:
>
> 10971 node = OMP_CLAUSE_CHAIN (node);
> 10972 if (node && omp_map_clause_descriptor_p (node))
> 10973 {
> 10974 gcc_assert (node != grp->grp_end);
> 10975 node = OMP_CLAUSE_CHAIN (node);
> 10976 }
>
> as the node == grp->grp_end.
This is now fixed and the testcase included (map-subarray-15.f90).
> * * *
>
>>
>> To support additional layers in the FE would probably mean saving in a
>> vector all allocatables found while walking the chain of references.
>> Then loop through the vector and call gfc_map_array_descriptor for
>> each element. Hopefully that will just work with no changes in the ME.
>>
>> If so, I'll come up with a revised version of the patch.
>
> Fingers crossed!
>
Multiple levels of nested allocatable components are now supported as
well and the relevant testcase included (map-subarray-14.f90).
There is still one specific case that is not handled properly (no
regression though) as demonstrated by another test
(map-subarray-16.f90): mapping both a subcomponent and the parent struct
on the same target directive.
This revised patch also incorporates a few formatting fixes.
Thanks,
Hi PA,
Paul-Antoine Arras wrote:
> This is now fixed and the testcase included (map-subarray-15.f90).
...
> Multiple levels of nested allocatable components are now supported as
> well and the relevant testcase included (map-subarray-14.f90).
> There is still one specific case that is not handled properly (no
> regression though) as demonstrated by another test
> (map-subarray-16.f90): mapping both a subcomponent and the parent
> struct on the same target directive.
>
> This revised patch also incorporates a few formatting fixes.
Thanks!
Can you update PR120505 after the commits to state what remains to be done? - Or,
alternatively and possibly better, create a new PR – and add a refer to it in
PR120505. If you file the PR first, you could update the should-fail TODO of
map-subarray-16.f90 to the new PR before committing.
(But I would refer to the new PR only after 1/2 and 2/2 have been committed, i.e
as last comment in PR120505 - such that it stands out at the end.)
Regarding map-subarray-16.f90: That not only "Demonstrate[s] a bogus error"
(preexisting) but it is also a check that this code doesn't ICE (which it
did with in an older version of this patch.)
* * *
Commit log text:
> This patch ensures that all such intervening allocatables in a reference chain
> is properly mapped. For the above example, the frontend has to create the
'are properly …' (now that you changed it to 'all allocatables').
Otherwise: LGTM. Thanks for the patch and sorry for the slow review!
Tobias
PS: Likewise for this patch, I'd prefer if the AuthorDate of the commit could
be newer than Dec 9, 2025, esp. given that later changes happened. (I know that
CommitterDate also exists, but it is not shown by default and I find old
AuthorDates confusing, esp. if it does not match the latest change. But this
is not a showstopper.)
@@ -3531,6 +3531,162 @@ get_symbol_rooted_namelist (hash_map<gfc_symbol *,
return NULL;
}
+/* 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)
+{
+ 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)
+ {
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_SIZE (node) = size_int (0);
+ }
+ 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
+ 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 (mid_desc_p ? OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT : 0);
+
+ return false;
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false,
@@ -3544,6 +3700,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 +4802,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 +4813,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 +4972,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 +4992,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");
@@ -11511,9 +11511,13 @@ omp_mapped_by_containing_struct (hash_map<tree_operand_hash_no_se,
}
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;
}
+ next:
decl = wsdecl;
}
@@ -13390,6 +13394,35 @@ 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. */
+ 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_SIZE (grp->grp_end)
+ == size_int (OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT))
+ {
+ 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
@@ -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);
/* This is the first sorted node in the struct sibling list. Use it
to recalculate the correct bias to use.
new file mode 100644
@@ -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\]\) map\(to:chunk\.tiles \[pointer set, len: 64\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: [0-9]+\]\) } 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: 64\]\) map\(attach_detach:\(struct tile_type\[0:\] \* restrict\) chunk\.tiles\.data \[bias: [0-9]+\]\) } 1 "original" } }
+
+end
new file mode 100644
@@ -0,0 +1,47 @@
+! { 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.
+
+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: 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: 64\]\) map\(detach:chunk\.tiles\.data \[bias: [0-9]+\]\)} 1 "gimple" } }
+
+end
@@ -667,6 +667,7 @@ enum omp_clause_fallback_kind {
OMP_CLAUSE_FALLBACK_NULL
};
+#define OMP_CLAUSE_MAP_SIZE_NEEDS_ADJUSTMENT -1
/* memory-order-clause on OpenMP atomic/flush constructs or
argument of atomic_default_mem_order clause. */
new file mode 100644
@@ -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
new file mode 100644
@@ -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