[OpenMP] Adds omp_target_is_accessible routine

[nicebert]

Adds omp_target_is_accessible routine.
Refactors common code from omp_target_is_present to work for both routines.

[llvmbot]

@llvm/pr-subscribers-clang
@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-offload

Author: None (nicebert)

Changes

Adds omp_target_is_accessible routine.
Refactors common code from omp_target_is_present to work for both routines.

---

Full diff: https://github.com/llvm/llvm-project/pull/138294.diff

5 Files Affected:

• (modified) offload/include/omptarget.h (+1)
• (modified) offload/libomptarget/OpenMP/API.cpp (+15-21)
• (modified) offload/libomptarget/exports (+1)
• (modified) offload/libomptarget/omptarget.cpp (+25)
• (added) offload/test/mapping/is_accessible.cpp (+43)

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..a0a126004d3f9 100644
--- a/offload/libomptarget/OpenMP/API.cpp
+++ b/offload/libomptarget/OpenMP/API.cpp
@@ -39,6 +39,8 @@ 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,33 +170,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,
                              size_t DstOffset, size_t SrcOffset, int DstDevice,
                              int SrcDevice) {
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..8716b33ce068a 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -198,6 +198,31 @@ 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..daf38e7afaf76
--- /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 <stdio.h>
+#include <iostream>
+#include <omp.h>
+#include <assert.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;
+}

[github-actions]

✅ With the latest revision this PR passed the C/C++ code formatter.

[shiltian]

FWIW, #143058 seems like doing the same thing.

[nicebert]

after Monday's discussion in the Accelerator subcommittee call I'm re-working the implementation to what was discussed & create a pr to clarify the wording in the spec.

[nicebert]

PR is broken and touches >5000 files

I think you're right. I must've messed up somewhere while rebasing

[nicebert]

oh I think it fixed itself by using the update branch feature on the PR @arsenm :)

[nicebert]

@carlobertolli can you review this. I've reworked the patch using hsa_amd_pointer_info as we discussed.

@CatherineMoore FIY