[OpenMP] Adds omp_target_is_accessible routine #138294

nicebert · 2025-05-02T15:26:05Z

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

llvmbot · 2025-05-02T15:26:40Z

@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 · 2025-05-02T15:28:19Z

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

offload/libomptarget/OpenMP/API.cpp

offload/libomptarget/device.cpp

offload/libomptarget/OpenMP/API.cpp

offload/include/omptarget.h

offload/libomptarget/OpenMP/API.cpp

offload/libomptarget/exports

offload/libomptarget/OpenMP/API.cpp

offload/include/omptarget.h

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).

Co-authored-by: Shilei Tian <[email protected]>

shiltian · 2025-07-28T15:30:38Z

FWIW, #143058 seems like doing the same thing.

offload/libomptarget/OpenMP/API.cpp

nicebert · 2025-07-30T13:59:33Z

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.

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).

…lvm-project into feat/omp_target_is_accessible

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).

nicebert · 2025-09-17T12:54:55Z

PR is broken and touches >5000 files

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

nicebert · 2025-09-17T13:03:35Z

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

nicebert · 2025-09-17T13:26:04Z

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

@CatherineMoore FIY

offload/libomptarget/device.cpp

…lvm-project into feat/omp_target_is_accessible

The implemetation is allowed to return -1 for the host device number. To be complient with the spec both the device number needs to be checked against both -1 as well as the value returned by omp_get_initial_device.

…lvm-project into feat/omp_target_is_accessible

shiltian · 2025-09-24T14:04:25Z

offload/libomptarget/OpenMP/API.cpp

+     DeviceNum, DPxPTR(Ptr), Size);
+
+  if (!Ptr) {
+    DP("Call to omp_target_is_accessible with NULL ptr returning false\n");


Does spec say so?

It does. OMP Specs 6.0, page 607 "If ptr is NULL, the routine returns zero". Please check the specs before asking.

shiltian · 2025-09-24T14:05:33Z

offload/libomptarget/OpenMP/API.cpp

+  // 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());


I wonder whether this is a fatal message or simple just return false? What does the spec say when the device number is invalid?

I agree with @shiltian : let's return false if the device doesn't exist. A warning is also in order.

shiltian · 2025-09-24T14:06:11Z

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

+    hsa_status_t Status =
+        hsa_amd_pointer_info(Ptr, &Info, malloc, &Count, &Agents);
+
+    if (Status != HSA_STATUS_SUCCESS)


Can we log here what is the reason of failing?

I guess that the two semantically relevant errors here would be:
HSA_EXT_POINTER_TYPE_UNKNOWN
and
HSA_STATUS_ERROR_INVALID_ARGUMENT (which we should never get because we already checked whether ptr is null). The rest of the errors "out of resources, hsa not initialized" would have been caught earlier on, is my guess

shiltian · 2025-09-24T14:06:30Z

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

+    for (uint32_t i = 0; i < Count; i++)
+      if (Agents[i].handle == getAgent().handle)
+        return Info.sizeInBytes >= Size;


nit:

Suggested change

for (uint32_t i = 0; i < Count; i++)

if (Agents[i].handle == getAgent().handle)

return Info.sizeInBytes >= Size;

for (uint32_t i = 0; i < Count; i++) {

if (Agents[i].handle == getAgent().handle)

return Info.sizeInBytes >= Size;

}

shiltian · 2025-09-24T14:07:20Z

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

+  /// Returns true if the plugin can guarantee that the associated
+  /// storage is accessible
+  bool isAccessiblePtr(const void *Ptr, size_t Size);
+  virtual bool isAccessiblePtrImpl(const void *Ptr, size_t Size) {


If this impl function is only for internal use, can we make them private?

shiltian · 2025-09-24T14:08:22Z

offload/test/mapping/is_accessible.cpp

+// RUN: | %fcheck-generic -check-prefix=NO_USM
+
+// REQUIRES: unified_shared_memory
+// REQUIRES: amdgpu


I'd make this a XFAIL for nvptx instead of a requirement, since it always returns false.

carlobertolli

With the requested changes, this patch is, in my opinion, ready to land.

carlobertolli · 2025-09-26T15:19:38Z

offload/libomptarget/OpenMP/API.cpp

+     DeviceNum, DPxPTR(Ptr), Size);
+
+  if (!Ptr) {
+    DP("Call to omp_target_is_accessible with NULL ptr returning false\n");


It does. OMP Specs 6.0, page 607 "If ptr is NULL, the routine returns zero". Please check the specs before asking.

carlobertolli · 2025-09-26T15:19:43Z

offload/libomptarget/OpenMP/API.cpp

+  // 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());


I agree with @shiltian : let's return false if the device doesn't exist. A warning is also in order.

carlobertolli · 2025-09-26T15:19:48Z

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

+    hsa_status_t Status =
+        hsa_amd_pointer_info(Ptr, &Info, malloc, &Count, &Agents);
+
+    if (Status != HSA_STATUS_SUCCESS)


I guess that the two semantically relevant errors here would be:
HSA_EXT_POINTER_TYPE_UNKNOWN
and
HSA_STATUS_ERROR_INVALID_ARGUMENT (which we should never get because we already checked whether ptr is null). The rest of the errors "out of resources, hsa not initialized" would have been caught earlier on, is my guess

carlobertolli · 2025-09-26T15:21:02Z

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

+    // If the pointer is unknown to HSA it's assumed a host pointer
+    // in that case the device can access it on unified memory support is
+    // enabled
+    return IsXnackEnabled;


Just a nit: there are GPUs that do not have xnack, but still have the ability to access host memory. This behavior is fine for now, but we will have to revisit based on GPU (later, not in this PR).

Then it should be documented.

llvmbot added the offload label May 2, 2025

shiltian reviewed May 2, 2025

View reviewed changes

offload/libomptarget/OpenMP/API.cpp Outdated Show resolved Hide resolved

nicebert force-pushed the feat/omp_target_is_accessible branch from e3a5812 to 0a9bb0f Compare May 5, 2025 16:02

nicebert force-pushed the feat/omp_target_is_accessible branch from 0a9bb0f to 3c22b15 Compare July 14, 2025 13:09

shiltian reviewed Jul 14, 2025

View reviewed changes

nicebert force-pushed the feat/omp_target_is_accessible branch from 3c22b15 to 3c092a7 Compare July 24, 2025 13:27

llvmbot added the backend:AMDGPU label Jul 24, 2025

mjklemm reviewed Jul 24, 2025

View reviewed changes

offload/libomptarget/exports Outdated Show resolved Hide resolved

offload/libomptarget/OpenMP/API.cpp Outdated Show resolved Hide resolved

offload/include/omptarget.h Outdated Show resolved Hide resolved

[OpenMP] Adds omp_target_is_accessible routine

025d36e

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).

nicebert force-pushed the feat/omp_target_is_accessible branch from 3c092a7 to 025d36e Compare July 28, 2025 14:48

llvmbot added the clang Clang issues not falling into any other category label Jul 28, 2025

nicebert and others added 2 commits July 28, 2025 16:49

Update offload/libomptarget/OpenMP/API.cpp

b33b27e

Co-authored-by: Shilei Tian <[email protected]>

Fix comment spelling

bf01578

Co-authored-by: Shilei Tian <[email protected]>

shiltian mentioned this pull request Jul 28, 2025

[OpenMp] Prototype OpenMP 5.1's omp_target_is_accessible #143058

Open

jprotze reviewed Jul 28, 2025

View reviewed changes

offload/libomptarget/OpenMP/API.cpp Outdated Show resolved Hide resolved

jprotze reviewed Jul 28, 2025

View reviewed changes

offload/libomptarget/OpenMP/API.cpp Outdated Show resolved Hide resolved

nicebert and others added 7 commits July 31, 2025 10:08

[OpenMP] Adds omp_target_is_accessible routine

d20f4d5

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).

[OpenMP] Adds omp_target_is_accessible routine

cb87242

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).

Merge branch 'feat/omp_target_is_accessible' of github.com:nicebert/l…

03d45cc

…lvm-project into feat/omp_target_is_accessible

Merge branch 'main' into feat/omp_target_is_accessible

21b1d6a

[OpenMP] Reverts omp_get_initial_device changes

95ab6fe

Merge branch 'feat/omp_target_is_accessible' of github.com:nicebert/l…

2436211

…lvm-project into feat/omp_target_is_accessible

Merge branch 'main' into feat/omp_target_is_accessible

c92c94f

nicebert changed the title ~~[OpenMP] Adds omp_target_is_accessible routine~~ [WIP][OpenMP] Adds omp_target_is_accessible routine Aug 5, 2025

nicebert and others added 2 commits August 6, 2025 14:29

Merge branch 'main' into feat/omp_target_is_accessible

eeb7604

[OpenMP] Adds omp_target_is_accessible routine

9d97424

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).

jhuber6 removed request for lanza, Endilll, daniel-grumberg, yota9, cyndyishida, ayermolo, DeinAlptraum, andykaylor and yozhu September 17, 2025 12:54

Merge branch 'main' into feat/omp_target_is_accessible

4345232

shiltian reviewed Sep 17, 2025

View reviewed changes

offload/libomptarget/device.cpp Outdated Show resolved Hide resolved

nicebert added 8 commits September 19, 2025 03:49

Fixes formatting and comment issues.

4b51745

Merge branch 'feat/omp_target_is_accessible' of github.com:nicebert/l…

0ef2e79

…lvm-project into feat/omp_target_is_accessible

Merge branch 'main' into feat/omp_target_is_accessible

454df9e

[OpenMP] Fixes check for host device number

79dd36f

The implemetation is allowed to return -1 for the host device number. To be complient with the spec both the device number needs to be checked against both -1 as well as the value returned by omp_get_initial_device.

Merge branch 'feat/omp_target_is_accessible' of github.com:nicebert/l…

0ccae1d

…lvm-project into feat/omp_target_is_accessible

Fixes formatting

588c394

Merge branch 'main' into feat/omp_target_is_accessible

b751e75

Merge branch 'main' into feat/omp_target_is_accessible

584553e

nicebert requested a review from mjklemm September 19, 2025 15:18

shiltian reviewed Sep 24, 2025

View reviewed changes

carlobertolli approved these changes Sep 26, 2025

View reviewed changes

[OpenMP] Adds omp_target_is_accessible routine #138294

Are you sure you want to change the base?

[OpenMP] Adds omp_target_is_accessible routine #138294

Uh oh!

Conversation

nicebert commented May 2, 2025

Uh oh!

llvmbot commented May 2, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented May 2, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

shiltian commented Jul 28, 2025

Uh oh!

Uh oh!

Uh oh!

nicebert commented Jul 30, 2025

Uh oh!

nicebert commented Sep 17, 2025

Uh oh!

nicebert commented Sep 17, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

nicebert commented Sep 17, 2025

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

carlobertolli left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

llvmbot commented May 2, 2025 •

edited

Loading

github-actions bot commented May 2, 2025 •

edited

Loading

nicebert commented Sep 17, 2025 •

edited

Loading