[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 14:35:57 PST 2023


https://github.com/jdoerfert created https://github.com/llvm/llvm-project/pull/74532

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.

>From 2836f081e2e1b510de9f28d9442352e174b7cefb 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            |  8 ++++++-
 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, 48 insertions(+), 11 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..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);
+}



More information about the Openmp-commits mailing list