Added initial file for testing mapping of pointer with not matching data item

[andrewkallai]

frontier clang 18.0.0 [OMPVV_WARNING: test_target_map_device_pointer.c:58] _ompvv_isOffloadingOn: 1, _ompvv_isSharedEnv: 1, test is invalid summit gcc 13.2.1 [OMPVV_WARNING: test_target_map_device_pointer.c:57] _ompvv_isOffloadingOn: 1, _ompvv_isSharedEnv: 1, test is invalid nvc 22.11 line 67: warning: statement is unreachable [code_is_unreachable] return errors; [OMPVV_WARNING: test_target_map_device_pointer.c:58] _ompvv_isOffloadingOn: 1, _ompvv_isSharedEnv: 1, test is invalid

[andrewkallai]

Per @tob2 Admittedly, I do not complete understand what the current version tries to do. However, I do agree with the GCC and the HPE/Cray compiler — namely, that due to 'defaultmap(none:pointer)' using 'p' inside of the first target region is invalid as it does not appear in any data-sharing clause.

---

I understand the purpose of this pull request/testcase correctly, the idea is to test OpenMP 5.2's "If a matching mapped list item is not found, the pointer retains its original value as per the firstprivate semantics described in Section 5.4.4." (last paragraph in 5.2's "5.8.6 Pointer Initialization for Device Data Environments")

In OpenMP 5.1, the same section (there: 2.21.7.2) had: "If a matching mapped list item is not found, the assigned initial value of the pointer is NULL unless otherwise specified (see Section 2.5.1)."

---

How about something like the following? With OpenMP 5.1 (or e.g. current GCC with non-shared-memory offloading), p1 and p2 are NULL in the target region, while with OpenMP 5.2 the are supposed to keep their value (as tested below).

#include <stdlib.h>
#include <stdint.h>
#include <omp.h>

int main () {
  int A[3] = {1,2,3};
  int *p1, *p2, *p3;
  intptr_t ip;
  p1 = (int*) 0x12345;
  p2 = (int*) omp_target_alloc (sizeof(int) * 5, omp_get_default_device());
  p3 = &A[0];
  ip = (intptr_t) p2;
  #pragma omp target enter data map(to:A)
  #pragma omp target
  {
    /* The pointees aren't mapped. */
    /* OpenMP 5.1 -> NULL */
    /* OpenMP 5.2 -> same value as before the target region. */
    if ((intptr_t) p1 != 0x12345) abort ();
    if ((intptr_t) p2 != ip) abort ();
    /* Mapped => must point to the corresponding device storage of 'A' */
    if (p3[0] != 1 || p3[1] != 2 || p3[2] != 3) abort ();
  }
  omp_target_free (p2, omp_get_default_device());
  return 0;
}

[Edited – namely adding:]

Notes to the example above:

• Implicit default mapping (=defaultmap(default:pointer)) is required. That is similar but not the same as firstprivate (= defaultmap(firstprivate:pointer)). With firstprivate(p1,p2) the code above works but firstprivate(p3) will then fail (unless shared memory). → https://github.com/SOLLVE/sollve_vv/blob/master/tests/5.0/requires/test_requires_unified_address.c for a testcase of the latter. (This is an OpenMP 5.0 testcase.)
• The code above only works if sizeof(*void) is the same on host and device (maybe also if it is larger on the host); it also requires that omp_target_alloc returns a normal/bare pointer (should be the case if pointer size is the same).
  • That's extremely common and likely other things break if that assumption does not hold, but the OpenMP spec does not guarantee this unless omp required unified_address is used. But unified_address goes beyond as it requires that for any address it is clear that host and device addresses are separate, which is not required for the testcase above

[spophale]

@andrewkallai to create a PR for changing behavior of "_ompvv_isSharedEnv" macro before this test is merged.

[andrewkallai]

grep -r "SHARED_ENV" tests/ --exclude=.F90 tests/4.5/application_kernels/alpaka_complex_template.cpp: OMPVV_TEST_SHARED_ENVIRONMENT; tests/4.5/application_kernels/reduction_separated_directives.c: OMPVV_TEST_SHARED_ENVIRONMENT; tests/4.5/application_kernels/reduction_separated_directives.cpp: OMPVV_TEST_SHARED_ENVIRONMENT; tests/4.5/target_enter_data/test_target_enter_data_if.c: OMPVV_TEST_AND_SET_SHARED_ENVIRONMENT(isSharedMemory) tests/4.5/target_enter_exit_data/test_target_enter_exit_data_map_global_array.c: OMPVV_TEST_SHARED_ENVIRONMENT; tests/4.5/target_teams_distribute/test_target_teams_distribute.c: OMPVV_TEST_SHARED_ENVIRONMENT tests/4.5/target_teams_distribute/test_target_teams_distribute_dist_schedule.c: OMPVV_TEST_SHARED_ENVIRONMENT; tests/5.0/requires/test_requires_reverse_offload.c: OMPVV_TEST_AND_SET_SHARED_ENVIRONMENT(is_shared_env); tests/5.0/target/test_target_device.c: OMPVV_TEST_AND_SET_SHARED_ENVIRONMENT(is_shared_env); tests/5.1/target/test_target_defaultmap_present_scalar.c: OMPVV_TEST_AND_SET_SHARED_ENVIRONMENT(is_shared_env);