From mboxrd@z Thu Jan  1 00:00:00 1970
Return-Path: <gcc-bugzilla@gcc.gnu.org>
Received: by sourceware.org (Postfix, from userid 48)
	id 294A13858C30; Thu, 16 Nov 2023 20:59:07 +0000 (GMT)
DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 294A13858C30
DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gcc.gnu.org;
	s=default; t=1700168347;
	bh=Pg7ZW5ANXhd+89wRrdGPQpYhW6nCncWC12zxkBHDLLI=;
	h=From:To:Subject:Date:In-Reply-To:References:From;
	b=ijQohDbWOcLNCxsOkqWZ+D7rZ+prro0esjWjWfJeEmCiqtWlVSmwbHNLnGaSdlkXM
	 Io9a8BetArrYOUtnOigX/bwx10AiNy0mG0WeAy5L+PvwnJIiAnQLzGIHhptZDHgUsa
	 SWMBfSquH5CmC4qc1qXBJCrMSPnCW73BTQOZXYKY=
From: "tschwinge at gcc dot gnu.org" <gcc-bugzilla@gcc.gnu.org>
To: gcc-bugs@gcc.gnu.org
Subject: [Bug fortran/111661] [OpenACC] Detach+Attach of DT component gives
 libgomp: [0x405140,96] is not mapped when running 'acc update' on DT var
 itself
Date: Thu, 16 Nov 2023 20:59:06 +0000
X-Bugzilla-Reason: CC
X-Bugzilla-Type: changed
X-Bugzilla-Watch-Reason: None
X-Bugzilla-Product: gcc
X-Bugzilla-Component: fortran
X-Bugzilla-Version: 14.0
X-Bugzilla-Keywords: openacc
X-Bugzilla-Severity: normal
X-Bugzilla-Who: tschwinge at gcc dot gnu.org
X-Bugzilla-Status: UNCONFIRMED
X-Bugzilla-Resolution: 
X-Bugzilla-Priority: P3
X-Bugzilla-Assigned-To: unassigned at gcc dot gnu.org
X-Bugzilla-Target-Milestone: ---
X-Bugzilla-Flags: 
X-Bugzilla-Changed-Fields: attachments.created
Message-ID: <bug-111661-4-KxwDN97FxG@http.gcc.gnu.org/bugzilla/>
In-Reply-To: <bug-111661-4@http.gcc.gnu.org/bugzilla/>
References: <bug-111661-4@http.gcc.gnu.org/bugzilla/>
Content-Type: text/plain; charset="UTF-8"
Content-Transfer-Encoding: quoted-printable
X-Bugzilla-URL: http://gcc.gnu.org/bugzilla/
Auto-Submitted: auto-generated
MIME-Version: 1.0
List-Id: <gcc-bugs.sourceware.org>

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

--- Comment #4 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Created attachment 56608
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=3D56608&action=3Dedit
'pr111661.c'

Before getting the Fortran case to work, let's indeed first look into some
conceptually corresponding C code:

(In reply to Tobias Burnus from comment #2)
> Looking at the code more closely, the problem is:
>=20
>   #pragma omp target oacc_exit_data map(delete:tab.val.data [len: 88])
>=20
> this tries to 'delete' the array descriptor - but as tab.val.data is part=
 of
> 'tab', this deletes all of "tab".

..., and indeed the same appears to happen in C:

> Compare the C example:

I completed this into a functional code, as follows (and attached).

> struct t { int *a; int n; };
> void f() {
>   struct t s;

Here, first initialize 's':

    s.n =3D 10;
    s.a =3D __builtin_malloc(s.n * sizeof *s.a);

Now, before 's.a', we first need to establish 's' itself:

    #pragma acc enter data copyin(s)

>   #pragma acc enter data copyin(s.a[:s.n])

Then, let's do something observable, for example:

    #pragma acc serial present(s)
      {
        for (int i =3D 0; i < s.n; ++i)
          s.a[i] =3D i * i;
      }

To be able to observe the computations, instead of:

>   #pragma acc exit data delete(s.a[:s.n])

..., do:

    #pragma acc exit data copyout(s.a[:s.n]) //finalize

After this, we expect 's' still to be alive:

    if (!acc_is_present(&s, sizeof s))
      __builtin_abort();

>   // for completeness, not relevant here:
>   #pragma acc exit data detach(s.a)
>   #pragma acc exit data delete(s.a)

I don't understand what you're doing here; I commented out these two.

Instead, now get rid of 's':

    #pragma acc exit data delete(s)
    if (acc_is_present(&s, sizeof s))
      __builtin_abort();

Verify results, and clean up:

    for (int i =3D 0; i < s.n; ++i)
      if (s.a[i] !=3D i * i)
        __builtin_abort();

    __builtin_free(s.a);

> }

This works fine with 'finalize' commented out.  However, with 'finalize'
enabled, we see:

> Again, here a 'finalize' would force
> the reference counts to zero and, hence, also delete 's' and not only the
> pointee/pointer target *s.a / s.a[0:.n] but also the pointer 's.a' itself.

... this behavior.

I've never in detail looked into the 'struct' mapping stuff -- I suppose the
problem here is not "simply" that '&s =3D=3D &s.a', and that's confusing the
runtime?

> QUESTION: Is the current code for C (and Fortran) correct according to the
> OpenACC specification or not?

Per my -- quick, not in-depth -- first look, I'd say the code is correct, a=
nd
thus GCC's behavior is wrong.

> FOLLOW UP QUESTION: If GCC's result is incorrect, what should the compiler
> do instead?

It has to treat the outer 's' separate from the inner 's.a'.  (..., even if
they happen to have the same address -- in case that's relevant here).

How does corresponding OpenMP code behave?=