Skip to content

Commit

Permalink
[OpenMP][FIX] Ensure we allow shared libraries without kernels (#74532)
Browse files Browse the repository at this point in the history
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.
  • Loading branch information
jdoerfert committed Dec 5, 2023
1 parent 9fbcdfc commit 9f87509
Show file tree
Hide file tree
Showing 6 changed files with 51 additions and 12 deletions.
7 changes: 5 additions & 2 deletions openmp/libomptarget/include/device.h
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down
12 changes: 10 additions & 2 deletions openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}

Expand Down Expand Up @@ -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;

Expand Down
17 changes: 10 additions & 7 deletions openmp/libomptarget/src/device.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand All @@ -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,
Expand Down
2 changes: 1 addition & 1 deletion openmp/libomptarget/src/omptarget.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
3 changes: 3 additions & 0 deletions openmp/libomptarget/test/Inputs/declare_indirect_func.c
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@

int func() { return 42; }
#pragma omp declare target indirect to(func)
22 changes: 22 additions & 0 deletions openmp/libomptarget/test/offloading/shared_lib_fp_mapping.c
Original file line number Diff line number Diff line change
@@ -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);
}

0 comments on commit 9f87509

Please sign in to comment.