Author: Shilei Tian Date: 2023-01-11T22:05:42-05:00 New Revision: 6e18277a51187ce8e861cdf0ab1395235e5b83d4
URL: https://github.com/llvm/llvm-project/commit/6e18277a51187ce8e861cdf0ab1395235e5b83d4 DIFF: https://github.com/llvm/llvm-project/commit/6e18277a51187ce8e861cdf0ab1395235e5b83d4.diff LOG: [OpenMP] Implement `omp_get_mapped_ptr` This patch implements the function `omp_get_mapped_ptr`. Fix #59945. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D141545 Added: openmp/libomptarget/test/api/omp_get_mapped_ptr.c Modified: clang/docs/OpenMPSupport.rst openmp/libomptarget/src/api.cpp openmp/libomptarget/src/exports Removed: ################################################################################ diff --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst index dca1486f9967f..16cb50afa1a13 100644 --- a/clang/docs/OpenMPSupport.rst +++ b/clang/docs/OpenMPSupport.rst @@ -111,7 +111,7 @@ OpenMP 5.0 Implementation Details The following table provides a quick overview over various OpenMP 5.0 features and their implementation status. Please post on the -`Discourse forums (Runtimes - OpenMP category)`_ for more +`Discourse forums (Runtimes - OpenMP category)`_ for more information or if you want to help with the implementation. @@ -257,8 +257,8 @@ OpenMP 5.1 Implementation Details The following table provides a quick overview over various OpenMP 5.1 features and their implementation status, as defined in the technical report 8 (TR8). -Please post on the -`Discourse forums (Runtimes - OpenMP category)`_ for more +Please post on the +`Discourse forums (Runtimes - OpenMP category)`_ for more information or if you want to help with the implementation. @@ -283,7 +283,7 @@ implementation. +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | omp_target_is_accessible routine | :none:`unclaimed` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ -| device | omp_get_mapped_ptr routine | :none:`unclaimed` | | +| device | omp_get_mapped_ptr routine | :none:`done` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ | device | new async target memory copy routines | :none:`unclaimed` | | +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+ @@ -365,12 +365,12 @@ implementation. OpenMP Extensions ================= -The following table provides a quick overview over various OpenMP +The following table provides a quick overview over various OpenMP extensions and their implementation status. These extensions are not currently defined by any standard, so links to associated LLVM documentation are provided. As these extensions mature, they will be considered for standardization. Please post on the -`Discourse forums (Runtimes - OpenMP category)`_ to provide feedback. +`Discourse forums (Runtimes - OpenMP category)`_ to provide feedback. +------------------------------+-----------------------------------------------------------------------------------+--------------------------+--------------------------------------------------------+ |Category | Feature | Status | Reviews | diff --git a/openmp/libomptarget/src/api.cpp b/openmp/libomptarget/src/api.cpp index f408449f01134..f96a2be2146e8 100644 --- a/openmp/libomptarget/src/api.cpp +++ b/openmp/libomptarget/src/api.cpp @@ -18,6 +18,7 @@ #include <climits> #include <cstdlib> #include <cstring> +#include <mutex> EXTERN int omp_get_num_devices(void) { TIMESCOPE(); @@ -318,3 +319,52 @@ EXTERN int omp_target_disassociate_ptr(const void *HostPtr, int DeviceNum) { 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, returning Ptr " DPxMOD ".\n", + DeviceNum, DPxPTR(Ptr)); + return const_cast<void *>(Ptr); + } + + int DevicesSize = omp_get_initial_device(); + { + std::lock_guard<std::mutex> 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<void *>(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 index fe1c015ac20b5..94a3ccabf5804 100644 --- a/openmp/libomptarget/src/exports +++ b/openmp/libomptarget/src/exports @@ -31,6 +31,7 @@ VERS1.0 { __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 index 0000000000000..a8e11f912d668 --- /dev/null +++ b/openmp/libomptarget/test/api/omp_get_mapped_ptr.c @@ -0,0 +1,39 @@ +// RUN: %libomptarget-compile-and-run-generic + +#include <assert.h> +#include <omp.h> +#include <stdlib.h> + +#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; +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
