[OpenMP][FIX] Ensure we do not read outside the device image (#74669)
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
https://github.com/llvm/llvm-project/pull/74550 picked the wrong one.
Fixes: https://github.com/llvm/llvm-project/issues/74582
GitOrigin-RevId: 0ace6ee73a6b7047d16f23170c67cdf358f34c34
diff --git a/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp b/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
index a3d16d3..0a19148 100644
--- a/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
+++ b/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 @@
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 @@
"%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/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp b/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
index 912e3d2..3c7d1ca 100644
--- a/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
@@ -785,9 +785,14 @@
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 @@
}
// 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/libomptarget/test/offloading/barrier_fence.c b/libomptarget/test/offloading/barrier_fence.c
index a0b672f..5d10964 100644
--- a/libomptarget/test/offloading/barrier_fence.c
+++ b/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/libomptarget/test/offloading/bug74582.c b/libomptarget/test/offloading/bug74582.c
new file mode 100644
index 0000000..c6a283b
--- /dev/null
+++ b/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
+ {}
+}
