commit 63f960a316c868aa1ec43a2a55839b8dcc037ed4
author Joseph Huber <huberjn@outlook.com> Sat Feb 03 15:28:20 2024 -0600
committer Copybara-Service <copybara-worker@google.com> Sat Feb 03 13:35:07 2024 -0800
tree 4f046b5068173e480e4447a6f63a034553f32454
parent db64313065bdf5d5479162c96f9e9dc8a7782800

[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.

GitOrigin-RevId: 2333865546cb6d4cda7b511ed07b8cb66a0d4eab

libomptarget/include/device.h

diff --git a/libomptarget/include/device.h b/libomptarget/include/device.h
index 3b40de9..bd28297 100644
--- a/libomptarget/include/device.h
+++ b/libomptarget/include/device.h
@@ -43,8 +43,6 @@
   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;

libomptarget/src/omptarget.cpp

diff --git a/libomptarget/src/omptarget.cpp b/libomptarget/src/omptarget.cpp
index 5b85249..74e23b0 100644
--- a/libomptarget/src/omptarget.cpp
+++ b/libomptarget/src/omptarget.cpp
@@ -132,9 +132,6 @@
 
 /// 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 @@
   if (Rc != OFFLOAD_SUCCESS)
     return Rc;
 
-  Device.HasMappedGlobalData = true;
-
   static Int32Envar DumpOffloadEntries =
       Int32Envar("OMPTARGET_DUMP_OFFLOAD_ENTRIES", -1);
   if (DumpOffloadEntries.get() == DeviceId)

libomptarget/test/offloading/dynamic_module_load.c

diff --git a/libomptarget/test/offloading/dynamic_module_load.c b/libomptarget/test/offloading/dynamic_module_load.c
index 935d402..5393f33 100644
--- a/libomptarget/test/offloading/dynamic_module_load.c
+++ b/libomptarget/test/offloading/dynamic_module_load.c
@@ -12,6 +12,9 @@
 #include <dlfcn.h>
 #include <stdio.h>
 int main(int argc, char **argv) {
+#pragma omp target
+  ;
+
   void *Handle = dlopen(argv[1], RTLD_NOW);
   int (*Foo)(void);