[OpenMP] Adds omp_target_is_accessible routine

Adds implementation of omp_target_is_accessible routine with
5.1 behaviour, checking if a host pointer is acccessible from a device
without running on the device (from the host).

Author: nicebert

clang/docs/OpenMPSupport.rst

@@ -256,7 +256,7 @@ implementation.
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device                       | device-specific environment variables                        | :none:`unclaimed`        |                                                                       |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
-| device                       | omp_target_is_accessible routine                             | :none:`unclaimed`        |                                                                       |
+| device                       | omp_target_is_accessible routine                             | :part:`worked on`        | https://github.com/llvm/llvm-project/pull/138294                      |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device                       | omp_get_mapped_ptr routine                                   | :good:`done`             | D141545                                                               |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+

offload/include/device.h

@@ -152,6 +152,9 @@ struct DeviceTy {
   /// Ask the device whether the runtime should use auto zero-copy.
   bool useAutoZeroCopy();
 
+  /// Ask the device whether it supports unified memory.
+  bool supportsUnifiedMemory();
+
   /// Check if there are pending images for this device.
   bool hasPendingImages() const { return HasPendingImages; }
 

offload/include/omptarget.h

@@ -280,6 +280,7 @@ int omp_get_initial_device(void);
 void *omp_target_alloc(size_t Size, int DeviceNum);
 void omp_target_free(void *DevicePtr, int DeviceNum);
 int omp_target_is_present(const void *Ptr, int DeviceNum);
+int omp_target_is_accessible(const void *Ptr, size_t Size, int DeviceNum);
 int omp_target_memcpy(void *Dst, const void *Src, size_t Length,
                       size_t DstOffset, size_t SrcOffset, int DstDevice,
                       int SrcDevice);

offload/libomptarget/OpenMP/API.cpp

@@ -93,7 +93,8 @@ EXTERN int omp_get_device_num(void) {
 EXTERN int omp_get_initial_device(void) {
   TIMESCOPE();
   OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
-  int HostDevice = omp_get_num_devices();
+  int NumDevices = omp_get_num_devices();
+  int HostDevice = NumDevices == 0 ? -1 : NumDevices;
   DP("Call to omp_get_initial_device returning %d\n", HostDevice);
   return HostDevice;
 }
@@ -195,6 +196,48 @@ EXTERN int omp_target_is_present(const void *Ptr, int DeviceNum) {
   return Rc;
 }
 
+/// Check whether a pointer is accessible from a device.
+/// the functionality is available in OpenMP 5.1 and later
+/// OpenMP 5.1
+/// omp_target_is_accessible checks whether a host pointer is accessible from a
+/// device OpenMP 6.0 removes restriction on pointer, allowing any pointer
+/// interpreted as a pointer in the address space of the given device.
+EXTERN int omp_target_is_accessible(const void *Ptr, size_t Size,
+                                    int DeviceNum) {
+  TIMESCOPE();
+  OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
+  DP("Call to omp_target_is_accessible for device %d, address " DPxMOD
+     ", size %zu\n",
+     DeviceNum, DPxPTR(Ptr), Size);
+
+  if (!Ptr) {
+    DP("Call to omp_target_is_accessible with NULL ptr returning false\n");
+    return false;
+  }
+
+  if (DeviceNum == omp_get_initial_device()) {
+    DP("Call to omp_target_is_accessible on host, returning true\n");
+    return true;
+  }
+
+  // the device number must refer to a valid device
+  auto DeviceOrErr = PM->getDevice(DeviceNum);
+  if (!DeviceOrErr)
+    FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
+
+  // for OpenMP 5.1 the routine checks whether a host pointer is accessible from
+  // the device this requires for the device to support unified shared memory
+  if (DeviceOrErr->supportsUnifiedMemory()) {
+    DP("Device %d supports unified memory, returning true\n", DeviceNum);
+    return true;
+  }
+
+  // functionality to check whether a device pointer is accessible from a device
+  // (OpenMP 6.0) from the host might not be possible
+  DP("Device %d does not support unified memory, returning false\n", DeviceNum);
+  return false;
+}
+
 EXTERN int omp_target_memcpy(void *Dst, const void *Src, size_t Length,
                              size_t DstOffset, size_t SrcOffset, int DstDevice,
                              int SrcDevice) {

offload/libomptarget/device.cpp

@@ -281,3 +281,7 @@ bool DeviceTy::useAutoZeroCopy() {
     return false;
   return RTL->use_auto_zero_copy(RTLDeviceID);
 }
+
+bool DeviceTy::supportsUnifiedMemory() {
+  return RTL->supports_unified_memory(RTLDeviceID);
+}

offload/libomptarget/exports

@@ -42,6 +42,7 @@ VERS1.0 {
     omp_get_initial_device;
     omp_target_alloc;
     omp_target_free;
+    omp_target_is_accessible;
     omp_target_is_present;
     omp_target_memcpy;
     omp_target_memcpy_rect;

offload/plugins-nextgen/amdgpu/src/rtl.cpp

@@ -2821,6 +2821,8 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
     return ((IsAPU || OMPX_ApuMaps) && IsXnackEnabled);
   }
 
+  bool supportsUnifiedMemoryImpl() override { return IsXnackEnabled; }
+
   /// Getters and setters for stack and heap sizes.
   Error getDeviceStackSize(uint64_t &Value) override {
     Value = StackSize;

offload/plugins-nextgen/common/include/PluginInterface.h

@@ -1003,6 +1003,10 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
   bool useAutoZeroCopy();
   virtual bool useAutoZeroCopyImpl() { return false; }
 
+  /// Returns true if the device has unified memory capabilities
+  bool supportsUnifiedMemory();
+  virtual bool supportsUnifiedMemoryImpl() { return false; }
+
   /// Allocate and construct a kernel object.
   virtual Expected<GenericKernelTy &> constructKernel(const char *Name) = 0;
 
@@ -1402,6 +1406,9 @@ struct GenericPluginTy {
   /// Returns if the plugin can support automatic copy.
   int32_t use_auto_zero_copy(int32_t DeviceId);
 
+  /// Returns if the the device supports unified memory.
+  int32_t supports_unified_memory(int32_t DeviceId);
+
   /// Look up a global symbol in the given binary.
   int32_t get_global(__tgt_device_binary Binary, uint64_t Size,
                      const char *Name, void **DevicePtr);

offload/plugins-nextgen/common/src/PluginInterface.cpp

@@ -1629,6 +1629,10 @@ Error GenericDeviceTy::syncEvent(void *EventPtr) {
 
 bool GenericDeviceTy::useAutoZeroCopy() { return useAutoZeroCopyImpl(); }
 
+bool GenericDeviceTy::supportsUnifiedMemory() {
+  return supportsUnifiedMemoryImpl();
+}
+
 Error GenericPluginTy::init() {
   if (Initialized)
     return Plugin::success();
@@ -2181,6 +2185,10 @@ int32_t GenericPluginTy::use_auto_zero_copy(int32_t DeviceId) {
   return getDevice(DeviceId).useAutoZeroCopy();
 }
 
+int32_t GenericPluginTy::supports_unified_memory(int32_t DeviceId) {
+  return getDevice(DeviceId).supportsUnifiedMemory();
+}
+
 int32_t GenericPluginTy::get_global(__tgt_device_binary Binary, uint64_t Size,
                                     const char *Name, void **DevicePtr) {
   assert(Binary.handle && "Invalid device binary handle");

offload/test/mapping/is_accessible.cpp

@@ -0,0 +1,40 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: env HSA_XNACK=1 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// RUN: %libomptarget-compilexx-generic
+// RUN: env HSA_XNACK=0 %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic -check-prefix=NO_USM
+
+// REQUIRES: unified_shared_memory
+// REQUIRES: amdgpu
+
+// CHECK: SUCCESS
+// NO_USM: Not accessible
+
+#include <assert.h>
+#include <iostream>
+#include <omp.h>
+#include <stdio.h>
+
+int main() {
+  int n = 10000;
+  int *a = new int[n];
+  int err = 0;
+
+  // program must be executed with HSA_XNACK=1
+  if (!omp_target_is_accessible(a, n * sizeof(int), /*device_num=*/0))
+    printf("Not accessible\n");
+  else {
+#pragma omp target teams distribute parallel for
+    for (int i = 0; i < n; i++)
+      a[i] = i;
+
+    for (int i = 0; i < n; i++)
+      if (a[i] != i)
+        err++;
+  }
+
+  printf("%s\n", err == 0 ? "SUCCESS" : "FAIL");
+  return err;
+}