[OpenMP][Offload] Add FB_NULLIFY map-type for use_device_ptr(fb_nullify).

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
 ^^^^^^^^^^^^

llvm/include/llvm/Frontend/OpenMP/OMPConstants.h

@@ -252,6 +252,10 @@ enum class OpenMPOffloadMappingFlags : uint64_t {
   // Attach pointer and pointee, after processing all other maps.
   // Applicable to map-entering directives. Does not change ref-count.
   OMP_MAP_ATTACH = 0x4000,
+  // When a lookup fails, fall back to using null as the translated pointer,
+  // instead of preserving the original pointer's value. Currently only
+  // useful in conjunction with RETURN_PARAM.
+  OMP_MAP_FB_NULLIFY = 0x8000,
   /// Signal that the runtime library should use args as an array of
   /// descriptor_dim pointers and use args_size as dims. Used when we have
   /// non-contiguous list items in target update directive

offload/include/omptarget.h

@@ -80,6 +80,10 @@ enum tgt_map_type {
   // Attach pointer and pointee, after processing all other maps.
   // Applicable to map-entering directives. Does not change ref-count.
   OMP_TGT_MAPTYPE_ATTACH = 0x4000,
+  // When a lookup fails, fall back to using null as the translated pointer,
+  // instead of preserving the original pointer's value. Currently only
+  // useful in conjunction with RETURN_PARAM.
+  OMP_TGT_MAPTYPE_FB_NULLIFY = 0x8000,
   // descriptor for non-contiguous target-update
   OMP_TGT_MAPTYPE_NON_CONTIG = 0x100000000000,
   // member of struct, member given by [16 MSBs] - 1

offload/libomptarget/omptarget.cpp

@@ -675,9 +675,46 @@ 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.
+        //
+        // OpenMP 6.1's `fb_nullify` fallback behavior: when the FB_NULLIFY bit
+        // is set by the compiler, e.g. for `use/need_device_ptr(fb_nullify)`),
+        // return `nullptr - Delta` when lookup fails.
+        if (ArgTypes[I] & OMP_TGT_MAPTYPE_FB_NULLIFY) {
+          TgtPtrBase = reinterpret_cast<void *>(
+              reinterpret_cast<uintptr_t>(nullptr) - Delta);
+          DP("Returning offsetted null pointer " DPxMOD
+             " as fallback (lookup failed)\n",
+             DPxPTR(TgtPtrBase));
+        } else {
+          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.