diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp --- a/openmp/libomptarget/src/api.cpp +++ b/openmp/libomptarget/src/api.cpp @@ -18,6 +18,7 @@ #include #include #include +#include EXTERN int omp_get_num_devices(void) { TIMESCOPE(); @@ -318,3 +319,51 @@ DP("omp_target_disassociate_ptr returns %d\n", Rc); return Rc; } + +EXTERN void *omp_get_mapped_ptr(const void *Ptr, int DeviceNum) { + TIMESCOPE(); + DP("Call to omp_get_mapped_ptr with ptr " DPxMOD ", device_num %d.\n", + DPxPTR(Ptr), DeviceNum); + + if (!Ptr) { + REPORT("Call to omp_get_mapped_ptr with nullptr.\n"); + return nullptr; + } + + if (DeviceNum == omp_get_initial_device()) { + REPORT("Device %d is initial device. No mapped pointer.\n", DeviceNum); + return nullptr; + } + + int DevicesSize = omp_get_initial_device(); + { + std::lock_guard LG(PM->RTLsMtx); + DevicesSize = PM->Devices.size(); + } + if (DevicesSize <= DeviceNum) { + DP("DeviceNum %d is invalid, returning nullptr.\n", DeviceNum); + return nullptr; + } + + if (!deviceIsReady(DeviceNum)) { + REPORT("Device %d is not ready, returning nullptr.\n", DeviceNum); + return nullptr; + } + + bool IsLast = false; + bool IsHostPtr = false; + auto &Device = *PM->Devices[DeviceNum]; + TargetPointerResultTy TPR = + Device.getTgtPtrBegin(const_cast(Ptr), 1, IsLast, + /*UpdateRefCount=*/false, + /*UseHoldRefCount=*/false, IsHostPtr); + if (!TPR.isPresent()) { + DP("Ptr " DPxMOD "is not present on device %d, returning nullptr.\n", + DPxPTR(Ptr), DeviceNum); + return nullptr; + } + + DP("omp_get_mapped_ptr returns " DPxMOD ".\n", DPxPTR(TPR.TargetPointer)); + + return TPR.TargetPointer; +} diff --git a/openmp/libomptarget/src/exports b/openmp/libomptarget/src/exports --- a/openmp/libomptarget/src/exports +++ b/openmp/libomptarget/src/exports @@ -31,6 +31,7 @@ __tgt_push_mapper_component; __kmpc_push_target_tripcount; __kmpc_push_target_tripcount_mapper; + omp_get_mapped_ptr; omp_get_num_devices; omp_get_device_num; omp_get_initial_device; diff --git a/openmp/libomptarget/test/api/omp_get_mapped_ptr.c b/openmp/libomptarget/test/api/omp_get_mapped_ptr.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/api/omp_get_mapped_ptr.c @@ -0,0 +1,39 @@ +// RUN: %libomptarget-compile-and-run-generic + +#include +#include +#include + +#define N 1024 +#define OFFSET 16 + +int main(int argc, char *argv[]) { + int *host_data = (int *)malloc(sizeof(int) * N); + void *device_ptr = omp_get_mapped_ptr(host_data, 0); + + assert(device_ptr == NULL && "the pointer should not be mapped right now"); + +#pragma omp target enter data map(to: host_data[:N]) + + device_ptr = omp_get_mapped_ptr(host_data, 0); + + assert(device_ptr && "the pointer should be mapped now"); + + void *ptr = NULL; + +#pragma omp target map(from: ptr) + { ptr = host_data; } + + assert(ptr == device_ptr && "wrong pointer mapping"); + + device_ptr = omp_get_mapped_ptr(host_data + OFFSET, 0); + + assert(device_ptr && "the pointer with offset should be mapped"); + +#pragma omp target map(from: ptr) + { ptr = host_data + OFFSET; } + + assert(ptr == device_ptr && "wrong pointer mapping"); + + return 0; +}