[OpenMP] Adds omp_target_is_accessible routine by nicebert · Pull Request #138294 · llvm/llvm-project (original) (raw)

@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:🎯: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 +#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; +}