[Bug fortran/111661] [OpenACC] Detach+Attach of DT component gives libgomp: [0x405140,96] is not mapped when running 'acc update' on DT var itself

[tschwinge at gcc dot gnu.org]

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

--- Comment #4 from Thomas Schwinge <tschwinge at gcc dot gnu.org> ---
Created attachment 56608
  --> https://gcc.gnu.org/bugzilla/attachment.cgi?id=56608&action=edit
'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:
> 
>   #pragma omp target oacc_exit_data map(delete:tab.val.data [len: 88])
> 
> 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 = 10;
    s.a = __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 = 0; i < s.n; ++i)
          s.a[i] = 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 = 0; i < s.n; ++i)
      if (s.a[i] != 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 == &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, and
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?