[Bug libgomp/106643] New: [gfortran + OpenACC] Allocate in module causes refcount error

[Bug libgomp/106643] New: [gfortran + OpenACC] Allocate in module causes refcount error

[hberre3 at gatech dot edu]

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

            Bug ID: 106643
           Summary: [gfortran + OpenACC] Allocate in module causes
                    refcount error
           Product: gcc
           Version: 13.0
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: hberre3 at gatech dot edu
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

I built GCC 13 from the default branch with offloading support for the new AMD
MI 210 GPUs by following the documented instructions. I ran into the following
runtime error when running our offloaded code written in Fortran leveraging
OpenACC:

```
/nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1153:
goacc_enter_data_internal: Assertion `n->refcount != REFCOUNT_INFINITY &&
n->refcount != REFCOUNT_LINK' failed.
```

I was able to create a minimal reproducible example for it:

p_main.f90
```
program p_main

    use m_macron

    !
==========================================================================

    implicit none

    call s_vars_init()
    call s_vars_read()
    call s_macron_init()

end program p_main
```

m_macron.f90
```
module m_macron

    use m_vars

    implicit none

    real(kind(0d0)), allocatable, dimension(:) :: valls
    !$acc declare create(valls)

contains

    subroutine s_macron_init()

        integer :: i

        print*, "size=", size

        print*, "allocate(valls(1:size))"
        allocate(valls(1:size))

        print*, "acc enter data create(valls(1:size))"
        !$acc enter data create(valls(1:size))

        print*, "!$acc update device(valls(1:size))"
        valls(size) = size - 2
        !$acc update device(valls(1:size))

        print*, valls(1:size)

        print*, "acc exit data delete(valls)"
        !$acc exit data delete(valls)

    end subroutine s_macron_init

end module m_macron
```

m_vars.f90
```
module m_vars

    integer :: size

contains

    subroutine s_vars_init()

        size = -100

    end subroutine s_vars_init

    subroutine s_vars_read()

        ! Namelist of the global parameters which may be specified by user
        namelist /user_inputs/ size

        open (1, FILE=trim("in.inp"), FORM='formatted', ACTION='read',
STATUS='old')
        read (1, NML=user_inputs); close (1)

    end subroutine s_vars_read

end module m_vars
```

in.inp
```
&user_inputs
size = 10
&end/
```

In order to generate this error, I had to create and dynamically allocate the
array in another module. I initially wrote this in a single F90 file but the
executable ran as expected. The error gets produced when running:
```
!$acc enter data create(valls(1:size))
```

Here is the full output when running the executable with the `GOMP_DEBUG`
environment variable set:
```
GOACC_data_start: mapnum=3, hostaddrs=0x7fff23d2efd0, size=0x7fff23d2efb0,
kinds=0x603102
  GOACC_data_start: prepare mappings
  GOACC_data_start: mappings prepared
 size=          10
 allocate(valls(1:size))
 acc enter data create(valls(1:size))
main:
/nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1153:
goacc_enter_data_internal: Assertion `n->refcount != REFCOUNT_INFINITY &&
n->refcount != REFCOUNT_LINK' failed.

Program received signal SIGABRT: Process abort signal.

Backtrace for this error:
#0  0x7f4aa4fcdb1f in ???
#1  0x7f4aa4fcda9f in ???
#2  0x7f4aa4fa0e04 in ???
#3  0x7f4aa4fa0cd8 in ???
#4  0x7f4aa4fc63f5 in ???
#5  0x7f4aa57b79e5 in goacc_enter_data_internal
        at
/nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1153
#6  0x7f4aa57b79e5 in goacc_enter_exit_data_internal
        at
/nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1405
#7  0x7f4aa57b8aab in GOACC_enter_data
        at
/nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1478
#8  0x4013b7 in __m_macron_MOD_s_macron_init
        at /nethome/hberre3/bug/m_macron.f90:22
#9  0x40182a in p_main
        at /nethome/hberre3/bug/p_main.f90:11
#10  0x401866 in main
        at /nethome/hberre3/bug/p_main.f90:3
run.sh: line 17: 3052573 Aborted                 (core dumped) ./main
```

I am not sure if this error is from libgomp or another part of the GCC
codebase. I assume it is related to an issue with the scoping of the array, as
it should be available throughout the entire program, per my reading of the
OpenACC spec:
```
The associated region is the implicit region associated with the function,
subroutine, or program in which the directive appears. If the directive appears
in the declaration section of a Fortran module subprogram or in a C or C++
global scope, the associated region is the implicit region for the whole
program. 
```

Our main code currently doesn't call `!$acc enter data create` for dynamically
allocated arrays since it relies on NVIDIA (/PGI) hooking into the `allocate`
call on the CPU. I ran into the above error when converting our
allocation/deallocation routines.

Here is the output of `gfortran -v`:
```
[hberre3@8:instinct]:~ $ ~/tools/gcc/13/bin/gfortran -v
Using built-in specs.
COLLECT_GCC=/nethome/hberre3/tools/gcc/13/bin/gfortran
COLLECT_LTO_WRAPPER=/nethome/hberre3/tools/gcc/13/libexec/gcc/x86_64-pc-linux-gnu/13.0.0/lto-wrapper
OFFLOAD_TARGET_NAMES=amdgcn-amdhsa
Target: x86_64-pc-linux-gnu
Configured with: /nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/configure
--build=x86_64-pc-linux-gnu --host=x86_64-pc-linux-gnu
--target=x86_64-pc-linux-gnu
--enable-offload-targets=amdgcn-amdhsa=/nethome/hberre3/tools/gcc/13/amdgcn-amdhsa
--enable-languages=c,c++,fortran,lto --disable-multilib
--prefix=/nethome/hberre3/tools/gcc/13
Thread model: posix
Supported LTO compression algorithms: zlib zstd
gcc version 13.0.0 20220811 (experimental) (GCC)
```

The code is compiled with `/nethome/hberre3/tools/gcc/13/bin/gfortran -O0 -g
-fopenacc '-foffload-options=-lgfortran -lm'
-foffload-options=amdgcn-amdhsa=-march=gfx90a -fno-exceptions`.

This example runs with NVFortran on NVIDIA GPUs. Thank you for taking a look!

[Bug libgomp/106643] [gfortran + OpenACC] Allocate in module causes refcount error

[hberre3 at gatech dot edu]

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

--- Comment #1 from Henry Le Berre <hberre3 at gatech dot edu> ---
Here is a more minimal version:

m_mod.f90:
```fortran
module m_mod

    real(kind(0d0)),allocatable, dimension(:) :: arrs
    !$acc declare create(arrs)

contains

    subroutine s_mod_init()

        allocate(arrs(10))

        !$acc enter data create(arrs(10))

        !$acc update device(arrs(10))

        !$acc exit data delete(arrs)

    end subroutine s_mod_init

end module m_mod
```


p_main.f90:
```

program p_main

    use m_mod

    call s_mod_init()

end program p_main
```

Here is a C version that runs without issue:

m_mod.c:
```
#include <stdio.h>
#include <stdlib.h>

double* arrs;

#pragma acc declare create(arrs)

void m_mod_init() {
    arrs = (double*)malloc(10*sizeof(double));

    #pragma acc enter data create(arrs[0:9])

    #pragma acc update device(arrs[0:9])

    #pragma acc exit data delete(arrs)
}
```

p_main.c:
```
extern void m_mod_init();

int main() {
    m_mod_init(); 
}
```

[Bug libgomp/106643] [gfortran + OpenACC] Allocate in module causes refcount error

[cvs-commit at gcc dot gnu.org]

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

--- Comment #2 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:da8e0e1191c5512244a752b30dea0eba83e3d10c

commit r13-3616-gda8e0e1191c5512244a752b30dea0eba83e3d10c
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Thu Oct 27 21:52:07 2022 +0200

    Support OpenACC 'declare create' with Fortran allocatable arrays, part I
[PR106643]

            PR libgomp/106643
            libgomp/
            * oacc-mem.c (goacc_enter_data_internal): Support
            OpenACC 'declare create' with Fortran allocatable arrays, part I.
            *
testsuite/libgomp.oacc-fortran/declare-allocatable-1-directive.f90:
            New.
            *
testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
            New.

[Bug libgomp/106643] [gfortran + OpenACC] Allocate in module causes refcount error

[cvs-commit at gcc dot gnu.org]

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

--- Comment #3 from CVS Commits <cvs-commit at gcc dot gnu.org> ---
The master branch has been updated by Thomas Schwinge <tschwinge@gcc.gnu.org>:

https://gcc.gnu.org/g:f6ce1e77bbf5d3a096f52e674bfd7354c6537d10

commit r13-3617-gf6ce1e77bbf5d3a096f52e674bfd7354c6537d10
Author: Thomas Schwinge <thomas@codesourcery.com>
Date:   Fri Oct 28 15:06:45 2022 +0200

    Support OpenACC 'declare create' with Fortran allocatable arrays, part II
[PR106643, PR96668]

            PR libgomp/106643
            PR fortran/96668
            libgomp/
            * oacc-mem.c (goacc_enter_data_internal): Support
            OpenACC 'declare create' with Fortran allocatable arrays, part II.
            *
testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.f90:
            Adjust.
            * testsuite/libgomp.oacc-fortran/pr106643-1.f90: New.

[Bug libgomp/106643] [gfortran + OpenACC] Allocate in module causes refcount error

[tschwinge at gcc dot gnu.org]

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

Thomas Schwinge <tschwinge at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Assignee|unassigned at gcc dot gnu.org      |tschwinge at gcc dot gnu.org
         Resolution|---                         |FIXED
                 CC|                            |tschwinge at gcc dot gnu.org
             Status|UNCONFIRMED                 |RESOLVED

--- Comment #4 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
(In reply to Henry Le Berre from comment #0)
> ```
> /nethome/hberre3/USERSCRATCH/build-gcc-amdgpu//gcc/libgomp/oacc-mem.c:1153:
> goacc_enter_data_internal: Assertion `n->refcount != REFCOUNT_INFINITY &&
> n->refcount != REFCOUNT_LINK' failed.
> ```

ACK, and thanks for your detailed report.

> In order to generate this error, I had to create and dynamically allocate
> the array in another module. I initially wrote this in a single F90 file but
> the executable ran as expected.

Maybe the difference was that the '!$acc declare create(valls)' was in a
different scope?  Global scope (Fortran 'module' scope) is important here; then
"the associated region is the implicit region for the whole program" (as you
also had noted), which triggers this code path.

There are cases where it's relevant, but here, the separate module/main program
files are not necessary for demonstrating the issue.

On the other hand, it's helpful to include an OpenACC compute construct where
for 'valls' there is no explicit or implicit data clause due to "exposed
variable access" (OpenACC 3.2 term).  That means, in a separate 'subroutine'
accessing 'valls' in '!$acc declare create(valls)' instead of dummy argument. 
(... as I've implemented in the test case.)

So this is now fixed for GCC 13; not planning on backporting to current GCC
release branches.


> Our main code currently doesn't call `!$acc enter data create` for
> dynamically allocated arrays since it relies on NVIDIA (/PGI) hooking into
> the `allocate` call on the CPU. I ran into the above error when converting
> our allocation/deallocation routines.

ACK, GCC release branches and master branch are still missing support for
OpenACC "Changes from Version 2.0 to 2.5": "The 'declare create' directive with
a Fortran 'allocatable' has new behavior".  Preliminary support exists on the
devel/omp/gcc-12 branch (and earlier development branches), but needs to be
revised for upstream submission.

[Bug libgomp/106643] [gfortran + OpenACC] Allocate in module causes refcount error

[pinskia at gcc dot gnu.org]

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

Andrew Pinski <pinskia at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
   Target Milestone|---                         |13.0