public inbox for gcc-bugs@sourceware.org
help / color / mirror / Atom feed
From: "patrick.begou@univ-grenoble-alpes.fr" <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: Fri, 13 Oct 2023 14:46:00 +0000	[thread overview]
Message-ID: <bug-111661-4-WW5Iug8sfO@http.gcc.gnu.org/bugzilla/> (raw)
In-Reply-To: <bug-111661-4@http.gcc.gnu.org/bugzilla/>

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

--- Comment #3 from Patrick Bégou <patrick.begou@univ-grenoble-alpes.fr> ---
Hi Tobias,

thanks for this information.
- yes removing the "finalize" make this small test case working. In my 
mind it should only remove the allocated attribute from the GPU with 
setting the count to zero. Is it because the attribute is an allocatable 
and not a pointer ? Is it the same behaviour with a pointer as attribute ?

- unfortunately this modification doesn't make some significant progress 
with porting my large code (things are much more complex) but with Gnu 
compilers Gdb is working. So it is a big step for investigating. I have 
isolated the low level data management to test this module 
independently. It works with ftn and Nvidia but not with Gnu at this 
time. I have to investigate.

Patrick

Le 13/10/2023 à 10:45, burnus at gcc dot gnu.org a écrit :
> https://gcc.gnu.org/bugzilla/show_bug.cgi?id=111661
>
> --- Comment #2 from Tobias Burnus <burnus at gcc dot gnu.org> ---
> @Patrick: It seems to work fine without "finalize".
>
> Can you check whether the big program then works for you?
> Usually, one should be able to do proper ref counting such that a
>   'finalize' -> force setting refcounts to zero
> shouldn't be needed.
>
> * * *
>
> 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".
>
>
> Compare the C example:
>
> struct t { int *a; int n; };
> void f() {
>    struct t s;
>    #pragma acc enter data copyin(s.a[:s.n])
>    #pragma acc exit data delete(s.a[:s.n])
>    // for completeness, not relevant here:
>    #pragma acc exit data detach(s.a)
>    #pragma acc exit data delete(s.a)
> }
>
>
> GCC does:
>
>   #pragma omp target oacc_enter_data map(struct:s [len: 1]) \
>       map(alloc:s.a [len: 8]) map(to:*_4 [len: _3]) map(attach:s.a [bias: 0])
>
>   #pragma omp target oacc_exit_data map(release:s.a [len: 8]) \
>       map(release:*_8 [len: _7]) map(detach:s.a [bias: 0])
>
>   #pragma omp target oacc_exit_data map(detach:s.a [bias: 0])
>   #pragma omp target oacc_exit_data map(release:s.a [len: 8])
>
> which seems to be at least consistent. 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.
>
> (BTW: Same result since GCC 10; GCC 9 rejects that code.)
>
>   * * *
>
> QUESTION: Is the current code for C (and Fortran) correct according to the
> OpenACC specification or not?
>
> FOLLOW UP QUESTION: If GCC's result is incorrect, what should the compiler do
> instead?
> And if it is correct, the question is: why do both ftn and nvfortran work in
> the same way?
>

  parent reply	other threads:[~2023-10-13 14:46 UTC|newest]

Thread overview: 5+ messages / expand[flat|nested]  mbox.gz  Atom feed  top
2023-10-02 10:40 [Bug fortran/111661] New: [OpenACC] Detach+Attach of DT component gives libgomp: [0x405140,96] is not mapped burnus at gcc dot gnu.org
2023-10-04 13:53 ` [Bug fortran/111661] [OpenACC] Detach+Attach of DT component gives libgomp: [0x405140,96] is not mapped when running 'acc update' on DT var itself patrick.begou@univ-grenoble-alpes.fr
2023-10-13  8:45 ` burnus at gcc dot gnu.org
2023-10-13 14:46 ` patrick.begou@univ-grenoble-alpes.fr [this message]
2023-11-16 20:59 ` tschwinge at gcc dot gnu.org

Reply instructions:

You may reply publicly to this message via plain-text email
using any one of the following methods:

* Save the following mbox file, import it into your mail client,
  and reply-to-all from there: mbox

  Avoid top-posting and favor interleaved quoting:
  https://en.wikipedia.org/wiki/Posting_style#Interleaved_style

* Reply using the --to, --cc, and --in-reply-to
  switches of git-send-email(1):

  git send-email \
    --in-reply-to=bug-111661-4-WW5Iug8sfO@http.gcc.gnu.org/bugzilla/ \
    --to=gcc-bugzilla@gcc.gnu.org \
    --cc=gcc-bugs@gcc.gnu.org \
    /path/to/YOUR_REPLY

  https://kernel.org/pub/software/scm/git/docs/git-send-email.html

* If your mail client supports setting the In-Reply-To header
  via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).