[Bug fortran/111661] New: [OpenACC] Detach+Attach of DT component gives libgomp: [0x405140,96] is not mapped

[burnus at gcc dot gnu.org via Gcc-bugs]

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=111661

            Bug ID: 111661
           Summary: [OpenACC] Detach+Attach of DT component gives libgomp:
                    [0x405140,96] is not mapped
           Product: gcc
           Version: 14.0
            Status: UNCONFIRMED
          Keywords: openacc
          Severity: normal
          Priority: P3
         Component: fortran
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: patrick.be...@univ-grenoble-alpes.fr,
                    tschwinge at gcc dot gnu.org
  Target Milestone: ---

Reported by Patrick. The following snippet gives:

libgomp: [0x406180,96] is not mapped

with GCC (mainline) and reportedly likewise GCC 13.2 but working with the
Nvidia and Cray compiler (see comments in the snipped). [Same result with an
older OG13 build: 13.1.1 20230710 [OG13].]


The problem seems to be the 'acc update' command. Here,
  0x406180
is the host address of 'tab':
          406180- tab - host
         1847B20- tab%val - host
      1304140000- tab - device
      1304140200- tab%val - device
          406180- tab - host
         18FF8C0- tab%val - host

libgomp: [0x406180,96] is not mapped


In other words, the DETACH seems to not only DETACH 'tab%val' but also somehow
unmap 'tab'?

#pragma omp target oacc_exit_data map(delete:tab.val.data [len: 88])
map(delete:MEM <real(kind=8)[0:]> [(real(kind=8)[0:] *)_34] [len: _33])
map(to:tab.val [pointer set, len: 88]) map(force_detach:tab.val.data [bias: 0])
finalize

...
          #pragma omp target oacc_update map(force_to:tab [len: 96])



!----------cut-------------------
use iso_c_binding
    implicit none

    type r2tab
             double precision, dimension(:,:), allocatable :: val
             integer :: dim1
             integer :: dim2
     end type r2tab
       type(r2tab) :: tab
       integer :: i,j
       integer(c_intptr_t) :: iloc(2)

       !$acc enter data copyin(tab)
       tab%dim1 = 10
       tab%dim2 = 20
       allocate (tab%val(tab%dim1,tab%dim2))
print '(z16,a)', loc(tab), "- tab - host"
print '(z16,a)', loc(tab%val), "- tab%val - host"
       !$acc enter data copyin(tab%val)
!$acc serial copyout(iloc)
iloc(1) = loc(tab)
iloc(2) = loc(tab%val)
!$acc end serial
print '(z16,a)', iloc(1), "- tab - device"
print '(z16,a)', iloc(2), "- tab%val - device"


!...   
       !$acc exit data delete(tab%val) finalize
! Works as is with nvfortran and CCE ftn but gfortran 13 requires
! additionally:
!      !$acc exit data delete(tab) finalize

       tab%dim1=11
       tab%dim2=30
       deallocate(tab%val)
       allocate (tab%val(tab%dim1,tab%dim2))
print '(z16,a)', loc(tab), "- tab - host"
print '(z16,a)', loc(tab%val), "- tab%val - host"
! For nvfortran and ftn:
       !$acc update device(tab)
! gfortran with the change above requires instead:
!         !$acc enter data copyin(tab)

       !$acc enter data create(tab%val)
!$acc parallel loop
  do j = 1, tab%dim2
    do i = 1, tab%dim1
      tab%val(i,j) = j * 100 + i
    end do
  end do
!$acc end  parallel loop
!$acc exit data copyout(tab%val) finalize
!$acc exit data delete(tab)

  do j = 1, tab%dim2
    do i = 1, tab%dim1
      if (tab%val(i,j) /= j * 100 + i) error stop
    end do
  end do
end