https://gcc.gnu.org/bugzilla/show_bug.cgi?id=120194
--- Comment #1 from Tobias Burnus <burnus at gcc dot gnu.org> --- Without having looked at it in depth, I think part of the problem is that 'varX' and 'varY' are in static memory. Thus, by construction, there is a device and a host version. Solution: (a) Ensure the date is actually mapped: - #pragma omp target + #pragma omp target map(always, tofrom: varX, varY) note the 'always' (b) Try the following (b.1) -#pragma omp declare target struct typeX varX; class typeY varY; -#pragma omp end declare target +#pragma omp declare target link(varX, varY) Note the 'link' And then (b.2) + #pragma target enter data (to: varX, varY) + #pragma omp target I had planned to do (b.2) automatically for USM - to get them auto intialized to the host version - and add some -f...=... flag to automatically change 'declare target' to 'declare target link', where that flag is enabled by default at least for 'requires self_maps', possibly also for 'requires unified_shared_memory'.