[Openmp-commits] [openmp] [OpenMP][FIX] Ensure we allow shared libraries without kernels (PR #74532)
Johannes Doerfert via Openmp-commits
openmp-commits at lists.llvm.org
Tue Dec 5 15:07:05 PST 2023
https://github.com/jdoerfert updated https://github.com/llvm/llvm-project/pull/74532
>From cf720763e9a42ee738740b55c44039f17d6b7acc Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes at jdoerfert.de>
Date: Tue, 5 Dec 2023 14:24:09 -0800
Subject: [PATCH] [OpenMP][FIX] Ensure we allow shared libraries without
kernels
This fixes two bugs and adds a test for them:
- A shared library with declare target functions but without kernels
should not error out due to missing globals.
- Enabling LIBOMPTARGET_INFO=32 should not deadlock in the presence of
indirect declare targets.
---
openmp/libomptarget/include/device.h | 7 ++++--
.../common/src/PluginInterface.cpp | 12 ++++++++--
openmp/libomptarget/src/device.cpp | 17 ++++++++------
openmp/libomptarget/src/omptarget.cpp | 2 +-
.../test/Inputs/declare_indirect_func.c | 3 +++
.../test/offloading/shared_lib_fp_mapping.c | 22 +++++++++++++++++++
6 files changed, 51 insertions(+), 12 deletions(-)
create mode 100644 openmp/libomptarget/test/Inputs/declare_indirect_func.c
create mode 100644 openmp/libomptarget/test/offloading/shared_lib_fp_mapping.c
diff --git a/openmp/libomptarget/include/device.h b/openmp/libomptarget/include/device.h
index 5146fc1444b44..ae7e0e11d4204 100644
--- a/openmp/libomptarget/include/device.h
+++ b/openmp/libomptarget/include/device.h
@@ -170,11 +170,14 @@ struct DeviceTy {
// Copy data from host to device
int32_t submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
AsyncInfoTy &AsyncInfo,
- HostDataToTargetTy *Entry = nullptr);
+ HostDataToTargetTy *Entry = nullptr,
+ DeviceTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr);
// Copy data from device back to host
int32_t retrieveData(void *HstPtrBegin, void *TgtPtrBegin, int64_t Size,
AsyncInfoTy &AsyncInfo,
- HostDataToTargetTy *Entry = nullptr);
+ HostDataToTargetTy *Entry = nullptr,
+ DeviceTy::HDTTMapAccessorTy *HDTTMapPtr = nullptr);
+
// Copy data from current device to destination device directly
int32_t dataExchange(void *SrcPtr, DeviceTy &DstDev, void *DstPtr,
int64_t Size, AsyncInfoTy &AsyncInfo);
diff --git a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
index 5a3fd140f27a3..7ca874bdfbf73 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
@@ -789,8 +789,10 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
sizeof(DeviceMemoryPoolTrackingTy),
&ImageDeviceMemoryPoolTracking);
if (auto Err =
- GHandler.readGlobalFromDevice(*this, *Image, TrackerGlobal))
- return Err;
+ GHandler.readGlobalFromDevice(*this, *Image, TrackerGlobal)) {
+ consumeError(std::move(Err));
+ continue;
+ }
DeviceMemoryPoolTracking.combine(ImageDeviceMemoryPoolTracking);
}
@@ -975,6 +977,12 @@ Error GenericDeviceTy::setupDeviceMemoryPool(GenericPluginTy &Plugin,
sizeof(DeviceMemoryPoolTrackingTy),
&DeviceMemoryPoolTracking);
GenericGlobalHandlerTy &GHandler = Plugin.getGlobalHandler();
+ if (auto Err = GHandler.readGlobalFromImage(*this, Image, TrackerGlobal)) {
+ [[maybe_unused]] std::string ErrStr = toString(std::move(Err));
+ DP("Avoid the memory pool: %s.\n", ErrStr.c_str());
+ return Error::success();
+ }
+
if (auto Err = GHandler.writeGlobalToDevice(*this, Image, TrackerGlobal))
return Err;
diff --git a/openmp/libomptarget/src/device.cpp b/openmp/libomptarget/src/device.cpp
index ad9563e04def4..fdc6da7a19d26 100644
--- a/openmp/libomptarget/src/device.cpp
+++ b/openmp/libomptarget/src/device.cpp
@@ -609,13 +609,14 @@ static void printCopyInfo(int DeviceId, bool H2D, void *SrcPtrBegin,
// Submit data to device
int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
- AsyncInfoTy &AsyncInfo,
- HostDataToTargetTy *Entry) {
+ AsyncInfoTy &AsyncInfo, HostDataToTargetTy *Entry,
+ DeviceTy::HDTTMapAccessorTy *HDTTMapPtr) {
if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
- HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor(Entry);
+ HDTTMapAccessorTy HDTTMap =
+ HostDataToTargetMap.getExclusiveAccessor(!!Entry || !!HDTTMapPtr);
LookupResult LR;
if (!Entry) {
- LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
+ LR = lookupMapping(HDTTMapPtr ? *HDTTMapPtr : HDTTMap, HstPtrBegin, Size);
Entry = LR.TPR.getEntry();
}
printCopyInfo(DeviceID, /* H2D */ true, HstPtrBegin, TgtPtrBegin, Size,
@@ -638,12 +639,14 @@ int32_t DeviceTy::submitData(void *TgtPtrBegin, void *HstPtrBegin, int64_t Size,
// Retrieve data from device
int32_t DeviceTy::retrieveData(void *HstPtrBegin, void *TgtPtrBegin,
int64_t Size, AsyncInfoTy &AsyncInfo,
- HostDataToTargetTy *Entry) {
+ HostDataToTargetTy *Entry,
+ DeviceTy::HDTTMapAccessorTy *HDTTMapPtr) {
if (getInfoLevel() & OMP_INFOTYPE_DATA_TRANSFER) {
- HDTTMapAccessorTy HDTTMap = HostDataToTargetMap.getExclusiveAccessor(Entry);
+ HDTTMapAccessorTy HDTTMap =
+ HostDataToTargetMap.getExclusiveAccessor(!!Entry || !!HDTTMapPtr);
LookupResult LR;
if (!Entry) {
- LR = lookupMapping(HDTTMap, HstPtrBegin, Size);
+ LR = lookupMapping(HDTTMapPtr ? *HDTTMapPtr : HDTTMap, HstPtrBegin, Size);
Entry = LR.TPR.getEntry();
}
printCopyInfo(DeviceID, /* H2D */ false, TgtPtrBegin, HstPtrBegin, Size,
diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp
index a9e22236dca27..2edbadaa6e02c 100644
--- a/openmp/libomptarget/src/omptarget.cpp
+++ b/openmp/libomptarget/src/omptarget.cpp
@@ -225,7 +225,7 @@ static int initLibrary(DeviceTy &Device) {
AsyncInfoTy AsyncInfo(Device);
void *DevPtr;
Device.retrieveData(&DevPtr, CurrDeviceEntryAddr, sizeof(void *),
- AsyncInfo);
+ AsyncInfo, /* Entry */ nullptr, &HDTTMap);
if (AsyncInfo.synchronize() != OFFLOAD_SUCCESS)
return OFFLOAD_FAIL;
CurrDeviceEntryAddr = DevPtr;
diff --git a/openmp/libomptarget/test/Inputs/declare_indirect_func.c b/openmp/libomptarget/test/Inputs/declare_indirect_func.c
new file mode 100644
index 0000000000000..20ac6617649ad
--- /dev/null
+++ b/openmp/libomptarget/test/Inputs/declare_indirect_func.c
@@ -0,0 +1,3 @@
+
+int func() { return 42; }
+#pragma omp declare target indirect to(func)
diff --git a/openmp/libomptarget/test/offloading/shared_lib_fp_mapping.c b/openmp/libomptarget/test/offloading/shared_lib_fp_mapping.c
new file mode 100644
index 0000000000000..8bd08ac5255c9
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/shared_lib_fp_mapping.c
@@ -0,0 +1,22 @@
+// clang-format off
+// RUN: %clang-generic -fPIC -shared %S/../Inputs/declare_indirect_func.c -o %T/liba.so -fopenmp-version=51
+// RUN: %libomptarget-compile-generic -L %T -l a -o %t -fopenmp-version=51
+// RUN: env LIBOMPTARGET_INFO=32 LD_LIBRARY_PATH=%T:$LD_LIBRARY_PATH %t | %fcheck-generic
+// clang-format on
+
+#include <stdio.h>
+
+extern int func(); // Provided in liba.so, returns 42
+typedef int (*fp_t)();
+
+int main() {
+ int x = 0;
+ fp_t fp = &func;
+ printf("TARGET\n");
+#pragma omp target map(from : x)
+ x = fp();
+ // CHECK: Copying data from device to host, {{.*}} Size=8
+ // CHECK: Copying data from device to host, {{.*}} Size=4
+ // CHECK: 42
+ printf("%i\n", x);
+}
More information about the Openmp-commits
mailing list