[OpenMP][FIX] Ensure we do not read outside the device image

[jdoerfert]

Before we expected all symbols in the device image to be backed up with data that we could read. However, uninitialized values are not. We now check for this case and avoid reading random memory.

This also replaces the correct readGlobalFromImage call with a isSymbolInImage check after
#74550 picked the wrong one.

Fixes: #74582

[llvmbot]

@llvm/pr-subscribers-openmp

Author: Johannes Doerfert (jdoerfert)

Changes

Before we expected all symbols in the device image to be backed up with data that we could read. However, uninitialized values are not. We now check for this case and avoid reading random memory.

This also replaces the correct readGlobalFromImage call with a isSymbolInImage check after
#74550 picked the wrong one.

Fixes: #74582

---

Full diff: https://github.com/llvm/llvm-project/pull/74669.diff

4 Files Affected:

• (modified) openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp (+22-3)
• (modified) openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp (+13-8)
• (modified) openmp/libomptarget/test/offloading/barrier_fence.c (-2)
• (added) openmp/libomptarget/test/offloading/bug74582.c (+13)

diff --git a/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp b/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
index a3d16d3a5bcff..0a19148ca4ec6 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
@@ -16,6 +16,10 @@
 
 #include "Shared/Utils.h"
 
+#include "llvm/BinaryFormat/ELF.h"
+#include "llvm/Support/Error.h"
+
+#include <cstdint>
 #include <cstring>
 
 using namespace llvm;
@@ -53,9 +57,15 @@ Error GenericGlobalHandlerTy::getGlobalMetadataFromELF(
     const ELF64LE::Shdr &Section, GlobalTy &ImageGlobal) {
 
   // The global's address is computed as the image begin + the ELF section
-  // offset + the ELF symbol value.
-  ImageGlobal.setPtr(advanceVoidPtr(
-      Image.getStart(), Section.sh_offset - Section.sh_addr + Symbol.st_value));
+  // offset + the ELF symbol value except for NOBITS sections that, as the name
+  // suggests, have no bits in the image. We still record the size and use
+  // nullptr to indicate there is no location.
+  if (Section.sh_type == ELF::SHT_NOBITS)
+    ImageGlobal.setPtr(nullptr);
+  else
+    ImageGlobal.setPtr(
+        advanceVoidPtr(Image.getStart(),
+                       Section.sh_offset - Section.sh_addr + Symbol.st_value));
 
   // Set the global's size.
   ImageGlobal.setSize(Symbol.st_size);
@@ -170,12 +180,21 @@ Error GenericGlobalHandlerTy::readGlobalFromImage(GenericDeviceTy &Device,
                          "%u bytes in the ELF image but %u bytes on the host",
                          HostGlobal.getName().data(), ImageGlobal.getSize(),
                          HostGlobal.getSize());
+  if (ImageGlobal.getPtr() == nullptr)
+    return Plugin::error("Transfer impossible because global symbol '%s' has "
+                         "no representation in the image (NOBITS sections)",
+                         HostGlobal.getName().data());
 
   DP("Global symbol '%s' was found in the ELF image and %u bytes will copied "
      "from %p to %p.\n",
      HostGlobal.getName().data(), HostGlobal.getSize(), ImageGlobal.getPtr(),
      HostGlobal.getPtr());
 
+  assert(Image.getStart() <= ImageGlobal.getPtr() &&
+         advanceVoidPtr(ImageGlobal.getPtr(), ImageGlobal.getSize()) <
+             advanceVoidPtr(Image.getStart(), Image.getSize()) &&
+         "Attempting to read outside the image!");
+
   // Perform the copy from the image to the host memory.
   std::memcpy(HostGlobal.getPtr(), ImageGlobal.getPtr(), HostGlobal.getSize());
 
diff --git a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
index 912e3d2c479b3..3c7d1ca899878 100644
--- a/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
@@ -785,9 +785,14 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
     GenericGlobalHandlerTy &GHandler = Plugin.getGlobalHandler();
     for (auto *Image : LoadedImages) {
       DeviceMemoryPoolTrackingTy ImageDeviceMemoryPoolTracking = {0, 0, ~0U, 0};
-      if (!GHandler.isSymbolInImage(*this, *Image,
-                                    "__omp_rtl_device_memory_pool_tracker"))
+      GlobalTy TrackerGlobal("__omp_rtl_device_memory_pool_tracker",
+                             sizeof(DeviceMemoryPoolTrackingTy),
+                             &ImageDeviceMemoryPoolTracking);
+      if (auto Err =
+              GHandler.readGlobalFromDevice(*this, *Image, TrackerGlobal)) {
+        consumeError(std::move(Err));
         continue;
+      }
       DeviceMemoryPoolTracking.combine(ImageDeviceMemoryPoolTracking);
     }
 
@@ -968,16 +973,16 @@ Error GenericDeviceTy::setupDeviceMemoryPool(GenericPluginTy &Plugin,
   }
 
   // Create the metainfo of the device environment global.
-  GlobalTy TrackerGlobal("__omp_rtl_device_memory_pool_tracker",
-                         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());
+  if (!GHandler.isSymbolInImage(*this, Image,
+                                "__omp_rtl_device_memory_pool_tracker")) {
+    DP("Skip the memory pool as there is no tracker symbol in the image.");
     return Error::success();
   }
 
+  GlobalTy TrackerGlobal("__omp_rtl_device_memory_pool_tracker",
+                         sizeof(DeviceMemoryPoolTrackingTy),
+                         &DeviceMemoryPoolTracking);
   if (auto Err = GHandler.writeGlobalToDevice(*this, Image, TrackerGlobal))
     return Err;
 
diff --git a/openmp/libomptarget/test/offloading/barrier_fence.c b/openmp/libomptarget/test/offloading/barrier_fence.c
index a0b672fb1a84a..5d1096478ed9e 100644
--- a/openmp/libomptarget/test/offloading/barrier_fence.c
+++ b/openmp/libomptarget/test/offloading/barrier_fence.c
@@ -7,8 +7,6 @@
 // UNSUPPORTED: aarch64-unknown-linux-gnu-LTO
 // UNSUPPORTED: x86_64-pc-linux-gnu
 // UNSUPPORTED: x86_64-pc-linux-gnu-LTO
-// UNSUPPORTED: amdgcn-amd-amdhsa
-// UNSUPPORTED: amdgcn-amd-amdhsa-LTO
 
 #include <omp.h>
 #include <stdio.h>
diff --git a/openmp/libomptarget/test/offloading/bug74582.c b/openmp/libomptarget/test/offloading/bug74582.c
new file mode 100644
index 0000000000000..c6a283bb93691
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/bug74582.c
@@ -0,0 +1,13 @@
+// RUN: %libomptarget-compile-generic && %libomptarget-run-generic
+// RUN: %libomptarget-compileopt-generic && %libomptarget-run-generic
+
+// Verify we do not read bits in the image that are not there (nobits section).
+
+#pragma omp begin declare target
+char BigUninitializedBuffer[4096 * 64] __attribute__((loader_uninitialized));
+#pragma omp end declare target
+
+int main() {
+#pragma omp target
+  {}
+}

[jhuber6]

Not directly related, but could we remove this function from GenericGlobalHandlerTy and just make this a utility in Elf.cpp like getSymbolAddr that returns said void pointer?

[jdoerfert]

We could. I guess it makes sense. Feel free to do it as a follow up?

[jhuber6]

Sure, it makes more sense to avoid propagating ELF specific stuff from the generic interface.

[jhuber6]

Sure, it makes more sense to avoid propagating ELF specific stuff from the generic interface.