[llvm] [OpenMP] Adds omp_target_is_accessible routine (PR #138294)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 14 06:10:04 PDT 2025
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/138294
>From 3c22b154c0548af210a9ba60ddfcdd929a7e563d Mon Sep 17 00:00:00 2001
From: Nicole Aschenbrenner <nicole.aschenbrenner at amd.com>
Date: Fri, 2 May 2025 09:58:23 -0400
Subject: [PATCH] [WIP][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).
---
offload/include/device.h | 3 ++
offload/include/omptarget.h | 1 +
offload/libomptarget/OpenMP/API.cpp | 43 +++++++++++++++++++
offload/libomptarget/device.cpp | 4 ++
offload/libomptarget/exports | 1 +
.../common/include/PluginInterface.h | 7 +++
.../common/src/PluginInterface.cpp | 8 ++++
offload/test/mapping/is_accessible.cpp | 43 +++++++++++++++++++
8 files changed, 110 insertions(+)
create mode 100644 offload/test/mapping/is_accessible.cpp
diff --git a/offload/include/device.h b/offload/include/device.h
index f4b10abbaa3fd..c82d947de6891 100644
--- a/offload/include/device.h
+++ b/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; }
diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h
index 6971780c7bdb5..8af8c4f659b35 100644
--- a/offload/include/omptarget.h
+++ b/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);
diff --git a/offload/libomptarget/OpenMP/API.cpp b/offload/libomptarget/OpenMP/API.cpp
index 4576f9bd06121..33af42c572d33 100644
--- a/offload/libomptarget/OpenMP/API.cpp
+++ b/offload/libomptarget/OpenMP/API.cpp
@@ -195,6 +195,49 @@ 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 || Size == 0) {
+ DP("Call to omp_target_is_accessible with NULL ptr or size 0, 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) {
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index f88e30ae9e76b..54ba7c30a90f4 100644
--- a/offload/libomptarget/device.cpp
+++ b/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);
+}
\ No newline at end of file
diff --git a/offload/libomptarget/exports b/offload/libomptarget/exports
index 2406776c1fb5f..0b770a2f1980a 100644
--- a/offload/libomptarget/exports
+++ b/offload/libomptarget/exports
@@ -37,6 +37,7 @@ VERS1.0 {
__kmpc_push_target_tripcount_mapper;
ompx_dump_mapping_tables;
omp_get_mapped_ptr;
+ omp_target_is_accessible;
omp_get_num_devices;
omp_get_device_num;
omp_get_initial_device;
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 162b149ab483e..1b0019835baef 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -981,6 +981,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;
@@ -1380,6 +1384,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);
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 81b9d423e13d8..ce2ee7f9af0ea 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -1626,6 +1626,10 @@ Error GenericDeviceTy::syncEvent(void *EventPtr) {
bool GenericDeviceTy::useAutoZeroCopy() { return useAutoZeroCopyImpl(); }
+bool GenericDeviceTy::supportsUnifiedMemory() {
+ return supportsUnifiedMemoryImpl();
+}
+
Error GenericPluginTy::init() {
if (Initialized)
return Plugin::success();
@@ -2178,6 +2182,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");
diff --git a/offload/test/mapping/is_accessible.cpp b/offload/test/mapping/is_accessible.cpp
new file mode 100644
index 0000000000000..2e8613729a139
--- /dev/null
+++ b/offload/test/mapping/is_accessible.cpp
@@ -0,0 +1,43 @@
+// RUN: %libomptarget-compilexx-generic
+// RUN: %libomptarget-run-generic 2>&1 \
+// RUN: | %fcheck-generic
+
+// REQUIRES: unified_shared_memory
+
+#include <assert.h>
+#include <iostream>
+#include <omp.h>
+#include <stdio.h>
+
+// The runtime considers unified shared memory to be always present.
+#pragma omp requires unified_shared_memory
+
+int main() {
+ int size = 10;
+ int *x = (int *)malloc(size * sizeof(int));
+ const int dev_num = omp_get_default_device();
+
+ int is_accessible = omp_target_is_accessible(x, size * sizeof(int), dev_num);
+ int errors = 0;
+ int uses_shared_memory = 0;
+
+#pragma omp target map(to : uses_shared_memory)
+ uses_shared_memory = 1;
+
+ assert(uses_shared_memory != is_accessible);
+
+ if (is_accessible) {
+#pragma omp target firstprivate(x)
+ for (int i = 0; i < size; i++)
+ x[i] = i * 3;
+
+ for (int i = 0; i < size; i++)
+ errors += (x[i] == (i * 3) ? 1 : 0);
+ }
+
+ free(x);
+ // CHECK: x overwritten 0 times
+ printf("x overwritten %d times\n", errors);
+
+ return errors;
+}
More information about the llvm-commits
mailing list