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?
