[Bug libgomp/103276] New: [openacc] Trying to map already mapped data

[Bug libgomp/103276] New: [openacc] Trying to map already mapped data

[tetra2005 at gmail dot com]

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

            Bug ID: 103276
           Summary: [openacc] Trying to map already mapped data
           Product: gcc
           Version: og11 (devel/omp/gcc-11)
            Status: UNCONFIRMED
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: tetra2005 at gmail dot com
                CC: jakub at gcc dot gnu.org
  Target Milestone: ---

Created attachment 51811
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=51811&action=edit
Reprocase

I've ran into a duplicate mapping error with OpenACC which seems to be caused
by invalid handling of ENTER DATA pragma.

The pragma, when applied to a parameter of derived type, also maps some
spurious local variable to device memory. After the function exits, another one
is called and tries to map it's own local variables which causes aliasing
conflict due to overlapping memory ranges in libgomp splay tree structures:
  libgomp: Trying to map into device [0x7fffffffcf80..0x7fffffffcfd8) object
when [0x7fffffffcfc8..0x7fffffffcfd0) is already mapped

The assembly looks like
  leaq  8(%rsp), %rax   # POINTER TO LOCAL STACK VAR
  ...
  subq  $8, %rsp
  movl  $2, %esi        # mapnum
  movq  %rdi, 24(%rsp)  # hostaddrs[0] (OK)
  movl  $-1, %edi       # flags_m
  movq  %rax, 32(%rsp)  # hostaddrs[1] (POINTS TO LOCAL STACK VAR)
  xorl  %eax, %eax
  pushq $0
  leaq  32(%rsp), %rdx  # hostaddrs
  call  GOACC_enter_exit_data
Here hostaddrs[1] points to a spurious variable in current stack frame.

Gimple code seems to be correct
  voidD.27 copyin_simple (struct simple & restrict varD.3961)
  { 
    struct .omp_data_t.1D.3962 D.3965;
    ...
    # .MEM_4 = VDEF <.MEM_3>
    D.3965.varD.3964 = &varD.3961;
but expanded RTL is not:
  ;; .omp_data_arr.2.var = &var;

  (insn 8 7 9 (parallel [
            (set (reg:DI 84)
                (plus:DI (reg/f:DI 77 virtual-stack-vars)
                    (const_int -24 [0xffffffffffffffe8])))
            (clobber (reg:CC 17 flags))
        ]) "repro.f90":11:28 -1
     (nil))

  (insn 9 8 0 (set (mem/f/c:DI (plus:DI (reg/f:DI 77 virtual-stack-vars)
                (const_int -8 [0xfffffffffffffff8])) [6 D.3965.varD.3964+0 S8
A64])
        (reg:DI 84)) "repro.f90":11:28 -1
     (nil))                                                                     

Reproduced with attached code on devel/omp/gcc-11 branch (commit f85ed2296,
2021 Nov 10) with
  gfortran -O2 -ffree-form -ffree-line-length-none -fopenacc
-foffload-options=-march=gfx908 -fdump-tree-all-all -fdump-rtl-all-all -S
repro.f90

[Bug libgomp/103276] [openacc] Trying to map already mapped data

[burnus at gcc dot gnu.org]

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

Tobias Burnus <burnus at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
           Keywords|                            |openacc

--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> ---
I only have GCC mainline at hand - and there it did not fail when adding:
  use REPRO
  type(simple) :: var
  call COPYIN_SIMPLE(var)
  end

This might be by chance, it could be fixed in mainline - or your example is too
much simplified.

At least as is it cannot fail as there is no main program but just a module. I
think a complete example would help. Or you confirm that it will fail as shown
after adding the 4 lines above...

[Bug libgomp/103276] [openacc] Trying to map already mapped data

[ygribov at gcc dot gnu.org]

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

Yury Gribov <ygribov at gcc dot gnu.org> changed:

           What    |Removed                     |Added
----------------------------------------------------------------------------
                 CC|                            |ygribov at gcc dot gnu.org

--- Comment #2 from Yury Gribov <ygribov at gcc dot gnu.org> ---
This might be by chance, it could be fixed in mainline - or your example is too
much simplified.

> At least as is it cannot fail as there is no main program but just a module.

I see. But is a runtime crash really mandatory here? It requires a bit of luck
so that address ranges of local variable in copyin_simple overlaps local
variables in another call. And even if it does, the crash will not be very
stable across branches (because stack layouts change, etc.).

My main goal was to provide a minimal self-explanatory usecase - there is a
clear reference to local stack object passed to GOACC_enter_exit_data which is
missing in the original code...

[Bug libgomp/103276] [openacc] Trying to map already mapped data

[pinskia at gcc dot gnu.org]

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

--- Comment #3 from Andrew Pinski <pinskia at gcc dot gnu.org> ---
(In reply to Yuri Gribov from comment #0)
> Here hostaddrs[1] points to a spurious variable in current stack frame.
> 
> Gimple code seems to be correct
>   voidD.27 copyin_simple (struct simple & restrict varD.3961)
>   { 
>     struct .omp_data_t.1D.3962 D.3965;
>     ...
>     # .MEM_4 = VDEF <.MEM_3>
>     D.3965.varD.3964 = &varD.3961;

No the gimple is wrong.  &var is taking the address of the argument. It should
just be var here if you what the reference points to rather than the address of
the decl holding the pointer.

[Bug libgomp/103276] [openacc] Trying to map already mapped data

[ygribov at gcc dot gnu.org]

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

--- Comment #4 from Yury Gribov <ygribov at gcc dot gnu.org> ---
(In reply to Andrew Pinski from comment #3)
> No the gimple is wrong.  &var is taking the address of the argument. It
> should just be var here if you what the reference points to rather than the
> address of the decl holding the pointer.

Thank you, makes sense. In that case the problem occurs at OMP lowering:
  voidD.27 copyin_simple (struct simple & restrict varD.3961)
  {
    var.0_1 = varD.3961;
    {
      D.3965.D.3963 = var.0_1;
      D.3965.varD.3964 = &varD.3961;
      #pragma omp target oacc_enter_exit_data map(to:*var.0_1 [len: 8])
map(alloc:varD.3961 [pointer assign, bias: 0])

[Bug libgomp/103276] [openacc] Trying to map already mapped data

[ygribov at gcc dot gnu.org]

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

--- Comment #5 from Yury Gribov <ygribov at gcc dot gnu.org> ---
Created attachment 51823
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=51823&action=edit
Runnable reprocase

I've attached another reprocase which reproduces the error at runtime (but let
me reiterate that the original minirepro is easier to analyze). It can be
compiled with
  /home/y.gribov/install/gcc-master/usr/local/bin/gfortran repro.f90 -O2
-ffree-form -fopenacc -foffload-options=-march=gfx908
It crashes at runtime with
  libgomp: Trying to map into device [0x7fffffffd8d0..0x7fffffffd928) object
when [0x7fffffffd918..0x7fffffffd920) is already mapped

Problem reproduces both on devel/omp/gcc-11 (commit f85ed229, Nov 10) and
master (commit 0136f25a, Nov 10).