[OpenMP] Preserve the original address when use_device_ptr/addr lookup fails.

clang/docs/OpenMPSupport.rst

@@ -266,6 +266,8 @@ implementation.
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device | has_device_addr clause on target construct | :none:`unclaimed` | |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
+| device | use_device_ptr/addr preserve host address when lookup fails | :good:`done` | https://github.com/llvm/llvm-project/pull/169438 |
++------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device | iterators in map clause or motion clauses | :none:`done` | https://github.com/llvm/llvm-project/pull/159112 |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device | indirect clause on declare target directive | :part:`In Progress` | |

clang/docs/ReleaseNotes.rst

@@ -773,6 +773,8 @@ OpenMP Support
 - Added parsing and semantic analysis support for ``need_device_ptr`` modifier
  to accept an optional fallback argument (``fb_nullify`` or ``fb_preserve``)
  with OpenMP >= 61.
+- ``use_device_ptr`` and ``use_device_addr`` now preserve the original host
+ address when lookup fails.
 
 Improvements
 ^^^^^^^^^^^^

offload/libomptarget/omptarget.cpp

@@ -675,9 +675,39 @@ int targetDataBegin(ident_t *Loc, DeviceTy &Device, int32_t ArgNum,
  DataSize, DPxPTR(TgtPtrBegin), (TPR.Flags.IsNewEntry ? "" : " not"));
 
  if (ArgTypes[I] & OMP_TGT_MAPTYPE_RETURN_PARAM) {
- uintptr_t Delta = (uintptr_t)HstPtrBegin - (uintptr_t)HstPtrBase;
- void *TgtPtrBase = (void *)((uintptr_t)TgtPtrBegin - Delta);
- DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ uintptr_t Delta = reinterpret_cast<uintptr_t>(HstPtrBegin) -
+ reinterpret_cast<uintptr_t>(HstPtrBase);
+ void *TgtPtrBase;
+ if (TgtPtrBegin) {
+ // Lookup succeeded, return device pointer adjusted by delta
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<uintptr_t>(TgtPtrBegin) - Delta);
+ DP("Returning device pointer " DPxMOD "\n", DPxPTR(TgtPtrBase));
+ } else {
+ // Lookup failed. So we have to decide what to do based on the
+ // requested fallback behavior.
+ //
+ // Treat "preserve" as the default fallback behavior, since as per
+ // OpenMP 5.1, for use_device_ptr/addr, when there's no corresponding
+ // device pointer to translate into, it's the user's responsibility to
+ // ensure that the host address is device-accessible.
+ //
+ // OpenMP 5.1, sec 2.14.2, target data construct, p 188, l26-31:
+ // If a list item that appears in a use_device_ptr clause ... does not
+ // point to a mapped object, it must contain a valid device address for
+ // the target device, and the list item references are instead converted
+ // to references to a local device pointer that refers to this device
+ // address.
+ //
+ // TODO: Add a new map-type bit to support OpenMP 6.1's `fb_nullify`
+ // and set the result to `nullptr - Delta`. Note that `fb_nullify` is
+ // already the default for `need_device_ptr`, but clang/flang do not
+ // support its codegen yet.
+ TgtPtrBase = reinterpret_cast<void *>(
+ reinterpret_cast<uintptr_t>(HstPtrBegin) - Delta);
+ DP("Returning host pointer " DPxMOD " as fallback (lookup failed).\n",
+ DPxPTR(TgtPtrBase));
+ }
  ArgsBase[I] = TgtPtrBase;
  }
 

offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_fallback.c

@@ -7,8 +7,6 @@
 // list-item is device-accessible, even if it was not
 // previously mapped.
 
-// XFAIL: *
-
 #include <stdio.h>
 int h[10];
 int *ph = &h[0];

offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_not_existing.cpp

@@ -8,15 +8,6 @@
 // Test for various cases of use_device_addr on an array-section.
 // The corresponding data is not previously mapped.
 
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
 int g, h[10];
 int *ph = &h[0];
 
@@ -36,7 +27,7 @@ struct S {
  int *mapped_ptr_ph3 =
  (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
  printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
  }
 
 // (B) use_device_addr/map: different operands, same base-pointer.
@@ -58,7 +49,7 @@ struct S {
  int *mapped_ptr_ph3 =
  (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
  printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
  }
 
 // (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -80,8 +71,7 @@ struct S {
  int **mapped_ptr_paa02 =
  (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
  printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
- mapped_ptr_paa02 != original_paa02,
- &paa[0][2] == (int **)nullptr + 2);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
  }
 
 // (F) use_device_addr/map: different operands, same base-array.
@@ -110,7 +100,7 @@ struct S {
  }
 
  int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
 
 // (H) use_device_addr/map: different base-pointers.
 // No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -122,7 +112,7 @@ struct S {
  int **mapped_ptr_paa0 =
  (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
  printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
- mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
  }
 
 // (I) use_device_addr/map: one map with different, one with same base-ptr.

offload/test/mapping/use_device_addr/target_data_use_device_addr_arrsec_ref_not_existing.cpp

@@ -8,15 +8,6 @@
 // Test for various cases of use_device_addr on an array-section on a reference.
 // The corresponding data is not previously mapped.
 
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
 int g_ptee;
 int &g = g_ptee;
 
@@ -37,15 +28,13 @@ struct S {
  int **original_paa02 = &paa[0][2];
 
 // (A) No corresponding map, lookup should fail.
-// EXPECTED: A: 1 1 1
-// CHECK: A: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: A: 1 1 1
 #pragma omp target data use_device_addr(ph[3 : 4])
  {
  int *mapped_ptr_ph3 =
  (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
  printf("A: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
  }
 
 // (B) use_device_addr/map: different operands, same base-pointer.
@@ -63,15 +52,13 @@ struct S {
 
 // (C) use_device_addr/map: different base-pointers.
 // No corresponding storage, lookup should fail.
-// EXPECTED: C: 1 1 1
-// CHECK: C: 1 1 0
-// FIXME: ph is not being privatized in the region.
+// CHECK: C: 1 1 1
 #pragma omp target data map(ph) use_device_addr(ph[3 : 4])
  {
  int *mapped_ptr_ph3 =
  (int *)omp_get_mapped_ptr(original_ph3, omp_get_default_device());
  printf("C: %d %d %d\n", mapped_ptr_ph3 == nullptr,
- mapped_ptr_ph3 != original_ph3, &ph[3] == (int *)nullptr + 3);
+ mapped_ptr_ph3 != original_ph3, &ph[3] == original_ph3);
  }
 
 // (D) use_device_addr/map: one of two maps with matching base-pointer.
@@ -95,8 +82,7 @@ struct S {
  int **mapped_ptr_paa02 =
  (int **)omp_get_mapped_ptr(original_paa02, omp_get_default_device());
  printf("E: %d %d %d\n", mapped_ptr_paa02 == nullptr,
- mapped_ptr_paa02 != original_paa02,
- &paa[0][2] == (int **)nullptr + 2);
+ mapped_ptr_paa02 != original_paa02, &paa[0][2] == original_paa02);
  }
 
 // (F) use_device_addr/map: different operands, same base-array.
@@ -125,7 +111,7 @@ struct S {
  }
 
  int *original_paa020 = &paa[0][2][0];
- int **original_paa0 = (int **)&paa[0];
+ void *original_paa0 = &paa[0];
 
 // (H) use_device_addr/map: different base-pointers.
 // No corresponding storage for use_device_addr opnd, lookup should fail.
@@ -137,7 +123,7 @@ struct S {
  int **mapped_ptr_paa0 =
  (int **)omp_get_mapped_ptr(original_paa0, omp_get_default_device());
  printf("H: %d %d %d\n", mapped_ptr_paa020 != nullptr,
- mapped_ptr_paa0 == nullptr, &paa[0] == nullptr);
+ mapped_ptr_paa0 == nullptr, &paa[0] == original_paa0);
  }
 
 // (I) use_device_addr/map: one map with different, one with same base-ptr.

offload/test/mapping/use_device_addr/target_data_use_device_addr_var_fallback.c

@@ -7,8 +7,6 @@
 // list-item is device-accessible, even if it was not
 // previously mapped.
 
-// XFAIL: *
-
 #include <stdio.h>
 int x;
 

offload/test/mapping/use_device_addr/target_data_use_device_addr_var_not_existing.cpp

@@ -8,15 +8,6 @@
 // Test for various cases of use_device_addr on a variable (not a section).
 // The corresponding data is not previously mapped.
 
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
 int g, h[10];
 int *ph = &h[0];
 
@@ -38,7 +29,7 @@ struct S {
  void *mapped_ptr_g =
  omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
  printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
- mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
  }
 
 // (B) Lookup should succeed.
@@ -58,7 +49,7 @@ struct S {
  void *mapped_ptr_h =
  omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
  printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
- mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
  }
 
 // (D) Lookup should succeed.
@@ -78,7 +69,7 @@ struct S {
  void *mapped_ptr_ph =
  omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
  printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
  }
 
 // (F) Lookup should succeed.
@@ -99,7 +90,7 @@ struct S {
  void *mapped_ptr_ph =
  omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
  printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
  }
 
 // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -119,7 +110,7 @@ struct S {
  void *mapped_ptr_paa =
  omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
  printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
  }
 
 // (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -130,7 +121,7 @@ struct S {
  void *mapped_ptr_paa =
  omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
  printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
  }
 
 // (K) Lookup should succeed.

offload/test/mapping/use_device_addr/target_data_use_device_addr_var_ref_not_existing.cpp

@@ -8,15 +8,6 @@
 // Test for various cases of use_device_addr on a reference variable.
 // The corresponding data is not previously mapped.
 
-// Note that this tests for the current behavior wherein if a lookup fails,
-// the runtime returns nullptr, instead of the original host-address.
-// That was compatible with OpenMP 5.0, where it was a user error if
-// corresponding storage didn't exist, but with 5.1+, the runtime needs to
-// return the host address, as it needs to assume that the host-address is
-// device-accessible, as the user has guaranteed it.
-// Once the runtime returns the original host-address when the lookup fails, the
-// test will need to be updated.
-
 int g_ptee;
 int &g = g_ptee;
 
@@ -45,7 +36,7 @@ struct S {
  void *mapped_ptr_g =
  omp_get_mapped_ptr(original_addr_g, omp_get_default_device());
  printf("A: %d %d %d\n", mapped_ptr_g == nullptr,
- mapped_ptr_g != original_addr_g, (void *)&g == nullptr);
+ mapped_ptr_g != original_addr_g, &g == original_addr_g);
  }
 
 // (B) Lookup should succeed.
@@ -65,7 +56,7 @@ struct S {
  void *mapped_ptr_h =
  omp_get_mapped_ptr(original_addr_h, omp_get_default_device());
  printf("C: %d %d %d\n", mapped_ptr_h == nullptr,
- mapped_ptr_h != original_addr_h, (void *)&h == nullptr);
+ mapped_ptr_h != original_addr_h, &h == original_addr_h);
  }
 
 // (D) Lookup should succeed.
@@ -85,7 +76,7 @@ struct S {
  void *mapped_ptr_ph =
  omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
  printf("E: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
  }
 
 // (F) Lookup should succeed.
@@ -106,7 +97,7 @@ struct S {
  void *mapped_ptr_ph =
  omp_get_mapped_ptr(original_addr_ph, omp_get_default_device());
  printf("G: %d %d %d\n", mapped_ptr_ph == nullptr,
- mapped_ptr_ph != original_addr_ph, (void *)&ph == nullptr);
+ mapped_ptr_ph != original_addr_ph, &ph == original_addr_ph);
  }
 
 // (H) Maps both pointee and pointer. Lookup for pointer should succeed.
@@ -126,7 +117,7 @@ struct S {
  void *mapped_ptr_paa =
  omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
  printf("I: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
  }
 
 // (J) Maps pointee only, but use_device_addr operand is pointer.
@@ -137,7 +128,7 @@ struct S {
  void *mapped_ptr_paa =
  omp_get_mapped_ptr(original_addr_paa, omp_get_default_device());
  printf("J: %d %d %d\n", mapped_ptr_paa == nullptr,
- mapped_ptr_paa != original_addr_paa, (void *)&paa == nullptr);
+ mapped_ptr_paa != original_addr_paa, &paa == original_addr_paa);
  }
 
 // (K) Lookup should succeed.

offload/test/mapping/use_device_addr/target_wrong_use_device_addr.c

@@ -1,5 +1,5 @@
 // RUN: %libomptarget-compile-generic -fopenmp-version=51 -g
-// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-fail-generic 2>&1 \
+// RUN: env LIBOMPTARGET_INFO=64 %libomptarget-run-generic 2>&1 \
 // RUN: | %fcheck-generic
 
 // FIXME: Fails due to optimized debugging in 'ptxas'
@@ -20,7 +20,8 @@ int main() {
 // counterpart
 #pragma omp target data use_device_addr(x)
  {
- // CHECK-NOT: device addr=0x[[#%x,HOST_ADDR:]]
+ // Even when the lookup fails, x should retain its host address.
+ // CHECK: device addr=0x[[#HOST_ADDR]]
  fprintf(stderr, "device addr=%p\n", x);
  }
  }