[Openmp-commits] [openmp] 0ace6ee - [OpenMP][FIX] Ensure we do not read outside the device image (#74669)
via Openmp-commits
openmp-commits at lists.llvm.org
Wed Dec 6 14:58:00 PST 2023
Author: Johannes Doerfert
Date: 2023-12-06T14:57:57-08:00
New Revision: 0ace6ee73a6b7047d16f23170c67cdf358f34c34
URL: https://github.com/llvm/llvm-project/commit/0ace6ee73a6b7047d16f23170c67cdf358f34c34
DIFF: https://github.com/llvm/llvm-project/commit/0ace6ee73a6b7047d16f23170c67cdf358f34c34.diff
LOG: [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
Added:
openmp/libomptarget/test/offloading/bug74582.c
Modified:
openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
openmp/libomptarget/plugins-nextgen/common/src/PluginInterface.cpp
openmp/libomptarget/test/offloading/barrier_fence.c
Removed:
################################################################################
diff --git a/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp b/openmp/libomptarget/plugins-nextgen/common/src/GlobalHandler.cpp
index a3d16d3a5bcff2..0a19148ca4ec6b 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 912e3d2c479b30..3c7d1ca8998787 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 a0b672fb1a84a2..5d1096478ed9e9 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 00000000000000..c6a283bb936914
--- /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