Skip to content

[OpenMp] Prototype OpenMP 5.1's omp_target_is_accessible #143058

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 5 commits into
base: main
Choose a base branch
from
Open

[OpenMp] Prototype OpenMP 5.1's omp_target_is_accessible #143058

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
0cecb7a
[OpenACC] Prototype OpenMP 5.1's omp_target_is_accessible
Jun 6, 2025
d347030
Return true when the pointer is contained and unified memory is enabled.
Jul 28, 2025
3ff2412
Add a check for pinned memory
Aug 6, 2025
4bcf315
Merge branch 'main' into yq/openmp/target-accessible-feature
yuanm-qy Aug 7, 2025
7293c09
Merge branch 'main' into yq/openmp/target-accessible-feature
yuanm-qy Aug 8, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions offload/include/omptarget.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
57 changes: 57 additions & 0 deletions offload/libomptarget/OpenMP/API.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -683,3 +683,60 @@ EXTERN void *omp_get_mapped_ptr(const void *Ptr, int DeviceNum) {

return TPR.TargetPointer;
}

EXTERN int omp_target_is_accessible(const void *Ptr, size_t Size,
int DeviceNum) {
TIMESCOPE();
// OpenMP 5.1, sec. 3.8.4 "omp_target_is_accessible", p. 417, L21-22:
// "This routine returns true if the storage of size bytes starting at the
// address given by Ptr is accessible from device device_num. Otherwise, it
// returns false."
//
// The meaning of "accessible" for unified shared memory is established in
// OpenMP 5.1, sec. 2.5.1 "requires directive". More generally, the specified
// host memory is accessible if it can be accessed from the device either
// directly (because of unified shared memory or because DeviceNum is the
// value returned by omp_get_initial_device()) or indirectly (because it's
// mapped to the device).
DP("Call to omp_target_is_accessible for device %d and address " DPxMOD "\n",
DeviceNum, DPxPTR(Ptr));

// FIXME: Is this right?
//
// Null pointer is permitted:
//
// OpenMP 5.1, sec. 3.8.4 "omp_target_is_accessible", p. 417, L15:
// "The value of ptr must be a valid host pointer or NULL (or C_NULL_PTR, for
// Fortran)."
//
// However, I found no specification of behavior in this case.
// omp_target_is_present has the same problem and is implemented the same way.
// Should Size have any effect on the result when Ptr is NULL?
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;
}

auto DeviceOrErr = PM->getDevice(DeviceNum);
if (!DeviceOrErr)
FATAL_MESSAGE(DeviceNum, "%s", toString(DeviceOrErr.takeError()).c_str());

// TODO: How does the spec intend for the Size=0 case to be handled?
// Currently, for the case where arr[N:M] is mapped, we return true for any
// address within arr[0:N+M]. However, Size>1 returns true only for arr[N:M].
// This is based on the discussion so far at the time of this writing at
// <https://github.com/llvm/llvm-project/issues/54899>. If the behavior
// changes, keep comments for omp_get_accessible_buffer in omp.h.var in sync.
TargetPointerResultTy TPR =
DeviceOrErr->getMappingInfo().getTgtPtrBegin(const_cast<void *>(Ptr), Size,
/*UpdateRefCount=*/false,
/*UseHoldRefCount=*/false);
int Rc = (TPR.isContained() || TPR.isHostPointer());
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is not sufficient. With the current implementation, omp_target_is_accessible requires an existing mapping.

The function is supposed to provide answers before setting up data mappings:

int main(void){
  int a, *b;
  cudaMallocHost((void**)&b, 1000);
  printf("a is accessible: %i\n", omp_target_is_accessible(&a, 4, 1));
  printf("b[:5] is accessible: %i\n", omp_target_is_accessible(b, 20, 1));
}

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thank you for your suggestion. I will make further modifications

DP("Call to omp_target_is_accessible returns %d\n", Rc);
return Rc;
}
1 change: 1 addition & 0 deletions offload/libomptarget/exports
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,7 @@ VERS1.0 {
omp_target_alloc;
omp_target_free;
omp_target_is_present;
omp_target_is_accessible;
omp_target_memcpy;
omp_target_memcpy_rect;
omp_target_memcpy_async;
Expand Down
4 changes: 4 additions & 0 deletions openmp/runtime/src/kmp_ftn_os.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -114,6 +114,7 @@
#define FTN_TARGET_ALLOC omp_target_alloc
#define FTN_TARGET_FREE omp_target_free
#define FTN_TARGET_IS_PRESENT omp_target_is_present
#define FTN_TARGET_IS_ACCESSIBLE omp_target_is_accessible
#define FTN_TARGET_MEMCPY omp_target_memcpy
#define FTN_TARGET_MEMCPY_RECT omp_target_memcpy_rect
#define FTN_TARGET_MEMSET omp_target_memset
Expand Down Expand Up @@ -263,6 +264,7 @@
#define FTN_TARGET_ALLOC omp_target_alloc_
#define FTN_TARGET_FREE omp_target_free_
#define FTN_TARGET_IS_PRESENT omp_target_is_present_
#define FTN_TARGET_IS_ACCESSIBLE omp_target_is_accessible_
#define FTN_TARGET_MEMCPY omp_target_memcpy_
#define FTN_TARGET_MEMCPY_RECT omp_target_memcpy_rect_
#define FTN_TARGET_ASSOCIATE_PTR omp_target_associate_ptr_
Expand Down Expand Up @@ -412,6 +414,7 @@
#define FTN_TARGET_ALLOC OMP_TARGET_ALLOC
#define FTN_TARGET_FREE OMP_TARGET_FREE
#define FTN_TARGET_IS_PRESENT OMP_TARGET_IS_PRESENT
#define FTN_TARGET_IS_ACCESSIBLE OMP_TARGET_IS_ACCESSIBLE
#define FTN_TARGET_MEMCPY OMP_TARGET_MEMCPY
#define FTN_TARGET_MEMCPY_RECT OMP_TARGET_MEMCPY_RECT
#define FTN_TARGET_ASSOCIATE_PTR OMP_TARGET_ASSOCIATE_PTR
Expand Down Expand Up @@ -559,6 +562,7 @@
#define FTN_TARGET_ALLOC OMP_TARGET_ALLOC_
#define FTN_TARGET_FREE OMP_TARGET_FREE_
#define FTN_TARGET_IS_PRESENT OMP_TARGET_IS_PRESENT_
#define FTN_TARGET_IS_ACCESSIBLE OMP_TARGET_IS_ACCESSIBLE_
#define FTN_TARGET_MEMCPY OMP_TARGET_MEMCPY_
#define FTN_TARGET_MEMCPY_RECT OMP_TARGET_MEMCPY_RECT_
#define FTN_TARGET_ASSOCIATE_PTR OMP_TARGET_ASSOCIATE_PTR_
Expand Down