mirrors/llvm-project
0
0
Fork 0
mirror of https://github.com/llvm/llvm-project.git synced 2025-04-26 05:46:06 +00:00
Code Issues Projects Releases Wiki Activity

[Libomptarget] Fix data mapping on dynamic loads (#80559)

Browse Source 
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.
This commit is contained in:
Joseph Huber 2024-02-03 15:28:20 -06:00 committed by GitHub
parent 9d00c34132
commit 2333865546
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
3 changed files with 3 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
openmp/libomptarget/include/device.h
View File

@ -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;

5
openmp/libomptarget/src/omptarget.cpp
View File

@ -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)

3
openmp/libomptarget/test/offloading/dynamic_module_load.c
View File

@ -12,6 +12,9 @@ int foo() {
#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);