[Openmp-commits] [openmp] [OpenMP][FIX] Ensure we do not read outside the device image (PR #74669)
via Openmp-commits
openmp-commits at lists.llvm.org
Wed Dec 6 14:42:06 PST 2023
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-openmp
Author: Johannes Doerfert (jdoerfert)
<details>
<summary>Changes</summary>
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
---
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
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
+ {}
+}
``````````
</details>
https://github.com/llvm/llvm-project/pull/74669
More information about the Openmp-commits
mailing list