From 2333865546cb6d4cda7b511ed07b8cb66a0d4eab Mon Sep 17 00:00:00 2001 From: Joseph Huber Date: Sat, 3 Feb 2024 15:28:20 -0600 Subject: [PATCH] [Libomptarget] Fix data mapping on dynamic loads (#80559) Summary: The current logic tries to map target mapping tables to the current device. Right now it assumes that data is only mapped a single time per device. This is only true if we have a single instance of the runtime running on a single program. However, in the case of dynamic library loads or shared libraries, this may happen multiple times. Given a case of a simple dynamic library load which has its own target kernel instruction, the current logic had only the first call to `__tgt_target_kernel` to the data mapping for that device. Then, when the next dynamic library load got called, it would see that the global were already mapped for that device and skip registering its own entires, even though they were distinct. This resulted in none of the mappings being done and hitting an assertion. This patch simply gets rid of this per-device check. The check should instead be on the host offloading entries. We already have logic that calls `continue` if we already have entries for that pointer, so we can simply rely on that instead. --- openmp/libomptarget/include/device.h | 2 -- openmp/libomptarget/src/omptarget.cpp | 5 ----- openmp/libomptarget/test/offloading/dynamic_module_load.c | 3 +++ 3 files changed, 3 insertions(+), 7 deletions(-) diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h index 3b40de959533b..bd2829722bb32 100644 --- a/openmp/libomptarget/include/device.h +++ b/openmp/libomptarget/include/device.h @@ -43,8 +43,6 @@ struct DeviceTy { PluginAdaptorTy *RTL; int32_t RTLDeviceID; - bool HasMappedGlobalData = false; - DeviceTy(PluginAdaptorTy *RTL, int32_t DeviceID, int32_t RTLDeviceID); // DeviceTy is not copyable DeviceTy(const DeviceTy &D) = delete; diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp index 5b852495b7501..74e23b06340c0 100644 --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -132,9 +132,6 @@ static uint64_t getPartialStructRequiredAlignment(void *HstPtrBase) { /// Map global data and execute pending ctors static int initLibrary(DeviceTy &Device) { - if (Device.HasMappedGlobalData) - return OFFLOAD_SUCCESS; - /* * Map global data */ @@ -294,8 +291,6 @@ static int initLibrary(DeviceTy &Device) { if (Rc != OFFLOAD_SUCCESS) return Rc; - Device.HasMappedGlobalData = true; - static Int32Envar DumpOffloadEntries = Int32Envar("OMPTARGET_DUMP_OFFLOAD_ENTRIES", -1); if (DumpOffloadEntries.get() == DeviceId) diff --git a/openmp/libomptarget/test/offloading/dynamic_module_load.c b/openmp/libomptarget/test/offloading/dynamic_module_load.c index 935d402ef2be1..5393f33e84f85 100644 --- a/openmp/libomptarget/test/offloading/dynamic_module_load.c +++ b/openmp/libomptarget/test/offloading/dynamic_module_load.c @@ -12,6 +12,9 @@ int foo() { #include #include int main(int argc, char **argv) { +#pragma omp target + ; + void *Handle = dlopen(argv[1], RTLD_NOW); int (*Foo)(void);