[Openmp-commits] [openmp] [OpenMP][FIX] Ensure we do not read outside the device image (PR #74669)
Johannes Doerfert via Openmp-commits
openmp-commits at lists.llvm.org
Wed Dec 6 14:41:51 PST 2023
https://github.com/jdoerfert created https://github.com/llvm/llvm-project/pull/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
>From 6f8927a775ac58b494d672aa73f21a41304fd5b8 Mon Sep 17 00:00:00 2001
From: Johannes Doerfert <johannes at jdoerfert.de>
Date: Wed, 6 Dec 2023 14:14:03 -0800
Subject: [PATCH] [OpenMP][FIX] Ensure we do not read outside the device image
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
---
.../common/src/GlobalHandler.cpp | 25 ++++++++++++++++---
.../common/src/PluginInterface.cpp | 21 ++++++++++------
.../test/offloading/barrier_fence.c | 2 --
.../libomptarget/test/offloading/bug74582.c | 13 ++++++++++
4 files changed, 48 insertions(+), 13 deletions(-)
create mode 100644 openmp/libomptarget/test/offloading/bug74582.c
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
+ {}
+}
More information about the Openmp-commits
mailing list