[OpenMP][FIX] Ensure we allow shared libraries without kernels

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

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;
 

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,

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;

openmp/libomptarget/test/Inputs/declare_indirect_func.c

@@ -0,0 +1,3 @@
+
+int func() { return 42; }
+#pragma omp declare target indirect to(func)

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);
+}