[llvm] [OpenMP] Adds omp_target_is_accessible routine (PR #138294)
via llvm-commits
llvm-commits at lists.llvm.org
Mon May 5 09:02:14 PDT 2025
https://github.com/nicebert updated https://github.com/llvm/llvm-project/pull/138294
>From 0a9bb0fdbf108d48446325e8b10878d2cdf8a5d4 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] [OpenMP] Adds omp_target_is_accessible routine
Adds omp_target_is_accessible routine.
Refactors common code from omp_target_is_present to work for both routines.
---
offload/include/omptarget.h | 1 +
offload/libomptarget/OpenMP/API.cpp | 39 +++++++++++------------
offload/libomptarget/exports | 1 +
offload/libomptarget/omptarget.cpp | 24 ++++++++++++++
offload/test/mapping/is_accessible.cpp | 43 ++++++++++++++++++++++++++
5 files changed, 87 insertions(+), 21 deletions(-)
create mode 100644 offload/test/mapping/is_accessible.cpp
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..6ed82719a3abb 100644
--- a/offload/libomptarget/OpenMP/API.cpp
+++ b/offload/libomptarget/OpenMP/API.cpp
@@ -39,6 +39,9 @@ EXTERN void ompx_dump_mapping_tables() {
using namespace llvm::omp::target::ompt;
#endif
+int checkTargetAddressMapping(const void *Ptr, size_t Size, int DeviceNum,
+ const char *Name);
+
void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
const char *Name);
void targetFreeExplicit(void *DevicePtr, int DeviceNum, int Kind,
@@ -168,31 +171,25 @@ EXTERN int omp_target_is_present(const void *Ptr, int DeviceNum) {
DP("Call to omp_target_is_present for device %d and address " DPxMOD "\n",
DeviceNum, DPxPTR(Ptr));
- if (!Ptr) {
- DP("Call to omp_target_is_present with NULL ptr, returning false\n");
- return false;
- }
-
- if (DeviceNum == omp_get_initial_device()) {
- DP("Call to omp_target_is_present on host, returning true\n");
- return true;
- }
-
- auto DeviceOrErr = PM->getDevice(DeviceNum);
- if (!DeviceOrErr)
- FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
-
// omp_target_is_present tests whether a host pointer refers to storage that
// is mapped to a given device. However, due to the lack of the storage size,
// only check 1 byte. Cannot set size 0 which checks whether the pointer (zero
// length array) is mapped instead of the referred storage.
- TargetPointerResultTy TPR =
- DeviceOrErr->getMappingInfo().getTgtPtrBegin(const_cast<void *>(Ptr), 1,
- /*UpdateRefCount=*/false,
- /*UseHoldRefCount=*/false);
- int Rc = TPR.isPresent();
- DP("Call to omp_target_is_present returns %d\n", Rc);
- return Rc;
+ return checkTargetAddressMapping(Ptr, 1, DeviceNum, "omp_target_is_present");
+}
+
+EXTERN int omp_target_is_accessible(const void *Ptr, size_t Size,
+ int DeviceNum) {
+ OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
+ DP("Call to omp_target_is_accessible for device %d and address " DPxMOD
+ " with size %zu\n",
+ DeviceNum, DPxPTR(Ptr), Size);
+
+ // omp_target_is_accessible tests whether a host pointer refers to storage
+ // that is mapped to a given device and is accessible from the device. The
+ // storage size is provided.
+ return checkTargetAddressMapping(Ptr, Size, DeviceNum,
+ "omp_target_is_accessible");
}
EXTERN int omp_target_memcpy(void *Dst, const void *Src, size_t Length,
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/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 5b25d955dd320..d3c5427737b3a 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -198,6 +198,30 @@ static int32_t getParentIndex(int64_t Type) {
return ((Type & OMP_TGT_MAPTYPE_MEMBER_OF) >> 48) - 1;
}
+int checkTargetAddressMapping(const void *Ptr, size_t Size, int DeviceNum,
+ const char *Name) {
+ if (!Ptr) {
+ DP("Call to %s with NULL ptr, returning false\n", Name);
+ return false;
+ }
+
+ if (DeviceNum == omp_get_initial_device()) {
+ DP("Call to %s on host, returning true\n", Name);
+ return true;
+ }
+
+ auto DeviceOrErr = PM->getDevice(DeviceNum);
+ if (!DeviceOrErr)
+ FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());
+
+ TargetPointerResultTy TPR = DeviceOrErr->getMappingInfo().getTgtPtrBegin(
+ const_cast<void *>(Ptr), Size, false, false);
+
+ int Rc = TPR.isPresent();
+ DP("Call to %s returns %d\n", Name, Rc);
+ return Rc;
+}
+
void *targetAllocExplicit(size_t Size, int DeviceNum, int Kind,
const char *Name) {
DP("Call to %s for device %d requesting %zu bytes\n", Name, DeviceNum, Size);
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