[Bug libgomp/127186] New: [OpenMP] With GOMP_RUNTIME_USM, in 'target'/'target data' still map static global variables (declare target)

[burnus at gcc dot gnu.org via Gcc-bugs]

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

            Bug ID: 127186
           Summary: [OpenMP] With GOMP_RUNTIME_USM, in 'target'/'target
                    data' still map static global variables (declare
                    target)
           Product: gcc
           Version: 17.0
            Status: UNCONFIRMED
          Keywords: openmp, wrong-code
          Severity: normal
          Priority: P3
         Component: libgomp
          Assignee: unassigned at gcc dot gnu.org
          Reporter: burnus at gcc dot gnu.org
                CC: jakub at gcc dot gnu.org, supers1ngular at gcc dot gnu.org
  Target Milestone: ---

With GOMP_RUNTIME_USM=enabled,
https://gcc.gnu.org/onlinedocs/libgomp/GOMP_005fRUNTIME_005fUSM.html , assume:

int result = 0;
#pragma omp declare target enter(result)

void test() {
  result = 3;
}

int main() {
  #pragma target map(always, tofrom: result)
    test ();
  printf ("result: %d\n", result);
}

Here, 'result' needs to be copied to the device and back - as the both the host
and the device have their own local 'result' variable.

For all other variables, mappings can be ignored (self map) as then directly
the host variable is accessed.

But global static variables exist both on the device and on the host and,
hence, those map clauses need to be still honored!

Longer version 'real-world' example:
   libgomp.c++/for-24.C
which uses such a results variable in a fancier way.

Note that both tests currently fail with a 'sorry' at runtime. [They only fail
when manually setting GOMP_RUNTIME_USM=enabled and a suitable device exists.]

* * *

Another testcase for this that sets GOMP_RUNTIME_USM=enabled
   testsuite/libgomp.c/usm_env_handling-1.c
but to avoid FAILS and skips unsupported bits tests
→ TODO: Enable those, after fixing this issue.

* * *

Commit r17-3807-g82104cde7ac0ba prepares for handling it by adding a SORRY
diagnostic for the failing cases to GOMP_target_ext and GOMP_target_data_ext.

Therefore: start there, when fixing it. For 'target' note that 'always' is not
enough (see below).

* * *

Note: While for 'target data', the issue only occurs with ALWAYS modifier, for
'target' it is also needed in general. Assume:

int A = 0;
#pragma omp declare target enter(A)

void f() {
  ++A;
}

int main()
{
#pragma omp target map(present, alloc: A)
  A = 5;

#pragma omp target
  f();

#pragma omp target update from(A)
__builtin_printf("A: %d\n");
}

This should print 'A: 6' - however, when self mapping is permitted and used
with an offload device due to GOMP_RUNTIME_USM=enabled:
In that case, in the first 'target' region, the A in 'A = 6' is using the
address of the host variable - but inside 'f()' the device variable - hence,
'A: 1' is currently printed.

Hence, for 'target' also the non-always case needs to be handled.

→ The second testcase currently does not fail with a 'sorry', i.e. yields wrong
code.