https://gcc.gnu.org/bugzilla/show_bug.cgi?id=93226
--- Comment #6 from Tobias Burnus <burnus at gcc dot gnu.org> ---
(In reply to Thomas Schwinge from comment #5)
> On OG15, for both nvptx and GCN offloading, I see:
...
The code has in the module:
integer :: D(N)
!$acc declare device_resident(D)
The problem is that on OG15, the compiler inserts at the beginning of the main
program:
#pragma omp target oacc_data map(force_alloc:d [len: 4096])
Thus, the variable is registered twice: Once through
GOMP_offload_register_ver
and then through
GOACC_data_start
leading to
libgomp: Trying to map into device [0x65d240..0x65e240) object
when [0x65d240..0x65e240) is already mapped
I bet that's a side-effect of commit
e361f9db605 Fortran "declare create"/allocate support for OpenACC
Maybe something like the following is needed?
(Light testing indicates that it fixes the issue.)
--- a/gcc/fortran/trans-openmp.cc
+++ b/gcc/fortran/trans-openmp.cc
@@ -4737,4 +4737,8 @@ gfc_trans_omp_clauses (stmtblock_t *block,
gfc_omp_clauses
prev = n;
+ if (cd == TOC_OPENACC_DECLARE
+ && n->sym->attr.oacc_declare_device_resident)
+ continue;
+
/* We do not want to include allocatable vars in a synthetic
"acc data" region created for "!$acc declare create" vars.