[OpenMP][FIX] Ensure we allow shared libraries without kernels #74532
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
|
@llvm/pr-subscribers-openmp Author: Johannes Doerfert (jdoerfert) ChangesThis fixes two bugs and adds a test for them:
Full diff: https://github.com/llvm/llvm-project/pull/74532.diff 6 Files Affected:
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..29c242448f813 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
@@ -790,7 +790,7 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
&ImageDeviceMemoryPoolTracking);
if (auto Err =
GHandler.readGlobalFromDevice(*this, *Image, TrackerGlobal))
- return Err;
+ continue;
DeviceMemoryPoolTracking.combine(ImageDeviceMemoryPoolTracking);
}
@@ -975,6 +975,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);
+}
|
jhuber6
reviewed
Dec 5, 2023
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.
jdoerfert
force-pushed
the
offload_prep11
branch
from
2836f08
to
cf72076
Compare
jhuber6
approved these changes
Dec 5, 2023
|
For reference: It appears that this caused #74582. I narrowed it down to a segmentation fault in the |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This fixes two bugs and adds a test for them: