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.

Reply via email to