[llvm] [offload] Add mechanism to return info to kernel replay tool (PR #192611)
Kevin Sala Penades via llvm-commits
llvm-commits at lists.llvm.org
Fri Apr 17 01:30:44 PDT 2026
https://github.com/kevinsala created https://github.com/llvm/llvm-project/pull/192611
This PR adds a mechanism to return information about a kernel replay to the kernel replay tool. This mechanism allows to verify the output of replays that use different configurations (e.g., different number of teams or threads) than the one used for recording.
It adds a new KernelExtraArgsTy structure that is only generated by the offload/libomptarget components (unlike KernelArgsTy). These runtime arguments may be used by future extensions that are not related to kernel record replay.
>From 97edf43f0992a18d654cbdb13c14872d861ab8e3 Mon Sep 17 00:00:00 2001
From: Kevin Sala <salapenades1 at llnl.gov>
Date: Fri, 17 Apr 2026 01:17:53 -0700
Subject: [PATCH] [offload] Add mechanism to return info to kernel replay tool
This PR adds a mechanism to return information about a kernel replay to
the kernel replay tool. This mechanism allows to verify the output of
replays that use different configurations (e.g., different number of
teams or threads) than the one used for recording.
It adds a new KernelExtraArgsTy structure that is only generated by the
offload/libomptarget components (unlike KernelArgsTy). These runtime
arguments may be used by future extensions that are not related to
kernel record replay.
---
offload/include/Shared/APITypes.h | 15 +++
offload/include/device.h | 1 +
offload/include/omptarget.h | 14 ++-
offload/liboffload/src/OffloadImpl.cpp | 2 +-
offload/libomptarget/device.cpp | 3 +-
offload/libomptarget/interface.cpp | 10 +-
offload/libomptarget/omptarget.cpp | 10 +-
offload/libomptarget/private.h | 3 +-
.../common/include/PluginInterface.h | 6 +-
.../common/include/RecordReplay.h | 45 ++++++---
.../common/src/PluginInterface.cpp | 14 ++-
.../common/src/RecordReplay.cpp | 95 +++++++++++++------
.../record-replay-diff-teams-threads.cpp | 46 +++++++++
.../record-replay-diff-threads.cpp | 46 +++++++++
.../kernelreplay/llvm-omp-kernel-replay.cpp | 10 +-
15 files changed, 249 insertions(+), 71 deletions(-)
create mode 100644 offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp
create mode 100644 offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp
diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h
index 6183686290bd4..40824596c3b9b 100644
--- a/offload/include/Shared/APITypes.h
+++ b/offload/include/Shared/APITypes.h
@@ -16,6 +16,7 @@
#include "Environment.h"
+#include "llvm/ADT/SmallString.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/Frontend/Offloading/Utility.h"
@@ -127,6 +128,20 @@ struct KernelLaunchParamsTy {
/// Ptrs to the Data entries. Only strictly required for the host plugin.
void **Ptrs = nullptr;
};
+
+/// The outcome of a kernel replay.
+struct KernelReplayOutcomeTy {
+ /// The path to the file that stores the output memory snapshot after the
+ /// kernel has been replayed.
+ llvm::SmallString<128> OutputFilepath;
+};
+
+/// Extra kernel arguments managed by the runtime components. Notice these
+/// arguments are additional to the ones in KernelArgsTy, which are usually
+/// generated by the compiler.
+struct KernelExtraArgsTy {
+ KernelReplayOutcomeTy *ReplayOutcome = nullptr;
+};
}
#endif // OMPTARGET_SHARED_API_TYPES_H
diff --git a/offload/include/device.h b/offload/include/device.h
index 06d21397c7377..af103c316c3cf 100644
--- a/offload/include/device.h
+++ b/offload/include/device.h
@@ -117,6 +117,7 @@ struct DeviceTy {
// Launch the kernel identified by \p TgtEntryPtr with the given arguments.
int32_t launchKernel(void *TgtEntryPtr, void **TgtVarsPtr,
ptrdiff_t *TgtOffsets, KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
AsyncInfoTy &AsyncInfo);
/// Synchronize device/queue/event based on \p AsyncInfo and return
diff --git a/offload/include/omptarget.h b/offload/include/omptarget.h
index c2233c5bfe2aa..0234e8fc55245 100644
--- a/offload/include/omptarget.h
+++ b/offload/include/omptarget.h
@@ -426,14 +426,12 @@ void __tgt_target_nowait_query(void **AsyncHandle);
/// Executes a target kernel by replaying recorded kernel arguments and
/// device memory.
-int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId, void *HostPtr,
- void *DeviceMemory, int64_t DeviceMemorySize,
- const llvm::offloading::EntryTy *Globals,
- int32_t NumGlobals, void **TgtArgs,
- ptrdiff_t *TgtOffsets, int32_t NumArgs,
- int32_t NumTeams, int32_t ThreadLimit,
- uint32_t SharedMemorySize,
- uint64_t LoopTripCount);
+int __tgt_target_kernel_replay(
+ ident_t *Loc, int64_t DeviceId, void *HostPtr, void *DeviceMemory,
+ int64_t DeviceMemorySize, const llvm::offloading::EntryTy *Globals,
+ int32_t NumGlobals, void **TgtArgs, ptrdiff_t *TgtOffsets, int32_t NumArgs,
+ int32_t NumTeams, int32_t ThreadLimit, uint32_t SharedMemorySize,
+ uint64_t LoopTripCount, KernelReplayOutcomeTy *ReplayOutcome);
void __tgt_set_info_flag(uint32_t);
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 77933e6291f4a..314788794842d 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1087,7 +1087,7 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
auto *KernelImpl = std::get<GenericKernelTy *>(Kernel->PluginImpl);
auto Err = KernelImpl->launch(*DeviceImpl, LaunchArgs.ArgPtrs, nullptr,
- LaunchArgs, AsyncInfoWrapper);
+ LaunchArgs, nullptr, AsyncInfoWrapper);
AsyncInfoWrapper.finalize(Err);
if (Err)
diff --git a/offload/libomptarget/device.cpp b/offload/libomptarget/device.cpp
index 277d4a428fa86..48aa29c3d4fbb 100644
--- a/offload/libomptarget/device.cpp
+++ b/offload/libomptarget/device.cpp
@@ -350,9 +350,10 @@ int32_t DeviceTy::notifyDataUnmapped(void *HstPtr) {
// Run region on device
int32_t DeviceTy::launchKernel(void *TgtEntryPtr, void **TgtVarsPtr,
ptrdiff_t *TgtOffsets, KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
AsyncInfoTy &AsyncInfo) {
return RTL->launch_kernel(RTLDeviceID, TgtEntryPtr, TgtVarsPtr, TgtOffsets,
- &KernelArgs, AsyncInfo);
+ &KernelArgs, KernelExtraArgs, AsyncInfo);
}
// Run region on device
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index 9f003ad67c2f4..9dd206d140c18 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -524,7 +524,7 @@ EXTERN int __tgt_target_kernel_replay(
int64_t DeviceMemorySize, const llvm::offloading::EntryTy *Globals,
int32_t NumGlobals, void **TgtArgs, ptrdiff_t *TgtOffsets, int32_t NumArgs,
int32_t NumTeams, int32_t ThreadLimit, uint32_t SharedMemorySize,
- uint64_t LoopTripCount) {
+ uint64_t LoopTripCount, KernelReplayOutcomeTy *ReplayOutcome) {
assert(PM && "Runtime not initialized");
OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
if (checkDevice(DeviceId, Loc)) {
@@ -541,10 +541,10 @@ EXTERN int __tgt_target_kernel_replay(
/*CodePtr=*/OMPT_GET_RETURN_ADDRESS);)
AsyncInfoTy AsyncInfo(*DeviceOrErr);
- int Rc =
- target_replay(Loc, *DeviceOrErr, HostPtr, DeviceMemory, DeviceMemorySize,
- Globals, NumGlobals, TgtArgs, TgtOffsets, NumArgs, NumTeams,
- ThreadLimit, SharedMemorySize, LoopTripCount, AsyncInfo);
+ int Rc = target_replay(
+ Loc, *DeviceOrErr, HostPtr, DeviceMemory, DeviceMemorySize, Globals,
+ NumGlobals, TgtArgs, TgtOffsets, NumArgs, NumTeams, ThreadLimit,
+ SharedMemorySize, LoopTripCount, AsyncInfo, ReplayOutcome);
if (Rc == OFFLOAD_SUCCESS)
Rc = AsyncInfo.synchronize();
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 4f7697639ff8e..f06654c639a8e 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -2350,7 +2350,7 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
#endif
Ret = Device.launchKernel(TgtEntryPtr, TgtArgs.data(), TgtOffsets.data(),
- KernelArgs, AsyncInfo);
+ KernelArgs, nullptr, AsyncInfo);
}
if (Ret != OFFLOAD_SUCCESS) {
@@ -2395,7 +2395,8 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
void **TgtArgs, ptrdiff_t *TgtOffsets, int32_t NumArgs,
int32_t NumTeams, int32_t ThreadLimit,
uint32_t SharedMemorySize, uint64_t LoopTripCount,
- AsyncInfoTy &AsyncInfo) {
+ AsyncInfoTy &AsyncInfo,
+ KernelReplayOutcomeTy *ReplayOutcome) {
int32_t DeviceId = Device.DeviceID;
int32_t NumSymbols = NumGlobals + 1;
@@ -2473,8 +2474,11 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
KernelArgs.ThreadLimit[2] = 1;
KernelArgs.DynCGroupMem = SharedMemorySize;
+ KernelExtraArgsTy KernelExtraArgs{};
+ KernelExtraArgs.ReplayOutcome = ReplayOutcome;
+
Ret = Device.launchKernel(Symbols[0].DevPtr, TgtArgs, TgtOffsets, KernelArgs,
- AsyncInfo);
+ &KernelExtraArgs, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Failed to launch kernel replay.";
return OFFLOAD_FAIL;
diff --git a/offload/libomptarget/private.h b/offload/libomptarget/private.h
index 7c2179d601543..31b295bda613e 100644
--- a/offload/libomptarget/private.h
+++ b/offload/libomptarget/private.h
@@ -37,7 +37,8 @@ extern int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
ptrdiff_t *TgtOffsets, int32_t NumArgs,
int32_t NumTeams, int32_t ThreadLimit,
uint32_t SharedMemorySize, uint64_t LoopTripCount,
- AsyncInfoTy &AsyncInfo);
+ AsyncInfoTy &AsyncInfo,
+ KernelReplayOutcomeTy *ReplayOutcome);
extern void handleTargetOutcome(bool Success, ident_t *Loc);
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 7cd0ebaae3b00..9e84d78834d9e 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -415,6 +415,7 @@ struct GenericKernelTy {
/// one used to initialize the kernel.
Error launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
AsyncInfoWrapperTy &AsyncInfoWrapper,
RecordReplayTy::HandleTy *RRHandle = nullptr) const;
virtual Error launchImpl(GenericDeviceTy &GenericDevice,
@@ -1019,7 +1020,9 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
/// Run the kernel associated with \p EntryPtr
Error launchKernel(void *EntryPtr, void **ArgPtrs, ptrdiff_t *ArgOffsets,
- KernelArgsTy &KernelArgs, __tgt_async_info *AsyncInfo);
+ KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
+ __tgt_async_info *AsyncInfo);
/// Initialize a __tgt_async_info structure.
Error initAsyncInfo(__tgt_async_info **AsyncInfoPtr);
@@ -1615,6 +1618,7 @@ struct GenericPluginTy {
/// Begin executing a kernel on the given device.
int32_t launch_kernel(int32_t DeviceId, void *TgtEntryPtr, void **TgtArgs,
ptrdiff_t *TgtOffsets, KernelArgsTy *KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
__tgt_async_info *AsyncInfoPtr);
/// Synchronize an asyncrhonous queue with the plugin runtime.
diff --git a/offload/plugins-nextgen/common/include/RecordReplay.h b/offload/plugins-nextgen/common/include/RecordReplay.h
index 2760e1037ed04..0929a533effa4 100644
--- a/offload/plugins-nextgen/common/include/RecordReplay.h
+++ b/offload/plugins-nextgen/common/include/RecordReplay.h
@@ -50,6 +50,15 @@ struct RecordReplayTy {
/// Describes the format of the recording and replaying.
enum FormatTy { Native = 0 };
+ /// Describes the file types that can be recorded.
+ enum FileTy {
+ PrologueSnapshot = 0,
+ EpilogueSnapshot,
+ Descriptor,
+ Globals,
+ Program
+ };
+
struct HandleTy {
const InstanceTy *Instance = nullptr;
bool Active = false;
@@ -103,11 +112,16 @@ struct RecordReplayTy {
size_t KernelHash = 0;
size_t LaunchConfigHash = 0;
+ /// When replaying, the entity requesting the replay may also need further
+ /// information about the the kernel's replay, such as the snapshot file.
+ KernelReplayOutcomeTy *ReplayOutcome = nullptr;
+
/// The number of occurrences during the execution.
mutable size_t Occurrences = 0;
InstanceTy(const GenericKernelTy &Kernel, uint32_t NumTeams,
- uint32_t NumThreads, uint32_t SharedMemorySize);
+ uint32_t NumThreads, uint32_t SharedMemorySize,
+ KernelReplayOutcomeTy *ReplayOutcome);
bool operator==(const InstanceTy &Other) const {
return (KernelHash == Other.KernelHash &&
@@ -170,6 +184,7 @@ struct RecordReplayTy {
/// instance is registered.
Expected<HandleTy>
recordPrologue(const GenericKernelTy &Kernel, const KernelArgsTy &KernelArgs,
+ const KernelExtraArgsTy *KernelExtraArgs,
const KernelLaunchParamsTy &LaunchParams, uint32_t NumTeams[3],
uint32_t NumThreads[3], uint32_t SharedMemorySize);
@@ -178,9 +193,9 @@ struct RecordReplayTy {
Error recordEpilogue(const GenericKernelTy &Kernel, HandleTy Handle);
/// Get a string with the filename.
- std::string getFilename(const InstanceTy &Instance, StringRef Suffix,
- bool IncludeDirectory = true) {
- return getFilenameImpl(Instance, Suffix, IncludeDirectory);
+ SmallString<128> getFilename(const InstanceTy &Instance, FileTy FileType,
+ bool IncludeDirectory = true) {
+ return getFilenameImpl(Instance, FileType, IncludeDirectory);
}
/// Allocates device memory from the record replay space.
@@ -192,7 +207,8 @@ struct RecordReplayTy {
/// as a new instance.
std::pair<const InstanceTy &, bool>
registerInstance(const GenericKernelTy &Kernel, uint32_t NumTeams,
- uint32_t NumThreads, uint32_t SharedMemorySize);
+ uint32_t NumThreads, uint32_t SharedMemorySize,
+ KernelReplayOutcomeTy *ReplayOutcome);
/// Record the prologue data.
virtual Error
@@ -211,9 +227,9 @@ struct RecordReplayTy {
const KernelLaunchParamsTy &LaunchParams) = 0;
/// Get a string with the filename.
- virtual std::string getFilenameImpl(const InstanceTy &Instance,
- StringRef Suffix,
- bool IncludeDirectory) = 0;
+ virtual SmallString<128> getFilenameImpl(const InstanceTy &Instance,
+ FileTy FileType,
+ bool IncludeDirectory) = 0;
};
/// The native kernel record replay support.
@@ -237,17 +253,20 @@ struct NativeRecordReplayTy : public RecordReplayTy {
const KernelLaunchParamsTy &LaunchParams) override;
/// Get a string with the filename.
- std::string getFilenameImpl(const InstanceTy &Instance, StringRef Suffix,
- bool IncludeDirectory) override;
+ SmallString<128> getFilenameImpl(const InstanceTy &Instance, FileTy FileType,
+ bool IncludeDirectory) override;
+
+ /// Get the extension for a recording file type.
+ StringRef getExtension(FileTy FileType);
/// Record a memory snapshot to a file.
- Error recordSnapshot(const std::string &Filename);
+ Error recordSnapshot(StringRef Filename);
/// Record the globals to a file.
- Error recordGlobals(const std::string &Filename);
+ Error recordGlobals(StringRef Filename);
/// Record the device image to a file.
- Error recordImage(const GenericKernelTy &Kernel, const std::string &Filename);
+ Error recordImage(const GenericKernelTy &Kernel, StringRef Filename);
};
} // namespace plugin
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 49b5edced9063..0d5cbf2e410aa 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -231,6 +231,7 @@ GenericKernelTy::prepareBlockMemory(GenericDeviceTy &GenericDevice,
Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
AsyncInfoWrapperTy &AsyncInfoWrapper,
RecordReplayTy::HandleTy *RRHandle) const {
llvm::SmallVector<void *, 16> Args;
@@ -278,9 +279,9 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
// blocks/threads.
RecordReplayTy *RecordReplay = GenericDevice.getRecordReplay();
if (RecordReplay) {
- auto RRHandleOrErr =
- RecordReplay->recordPrologue(*this, KernelArgs, LaunchParams, NumBlocks,
- NumThreads, DynBlockMemConf.NativeSize);
+ auto RRHandleOrErr = RecordReplay->recordPrologue(
+ *this, KernelArgs, KernelExtraArgs, LaunchParams, NumBlocks, NumThreads,
+ DynBlockMemConf.NativeSize);
if (!RRHandleOrErr)
return RRHandleOrErr.takeError();
if (RRHandle)
@@ -1105,6 +1106,7 @@ Error GenericDeviceTy::dataFill(void *TgtPtr, const void *PatternPtr,
Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
ptrdiff_t *ArgOffsets,
KernelArgsTy &KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
__tgt_async_info *AsyncInfo) {
AsyncInfoWrapperTy AsyncInfoWrapper(*this,
RecordReplay ? nullptr : AsyncInfo);
@@ -1126,7 +1128,7 @@ Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
RecordReplayTy::HandleTy RRHandle;
auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs,
- AsyncInfoWrapper, &RRHandle);
+ KernelExtraArgs, AsyncInfoWrapper, &RRHandle);
// 'finalize' here to guarantee next record-replay actions are in-sync
AsyncInfoWrapper.finalize(Err);
@@ -1654,9 +1656,11 @@ int32_t GenericPluginTy::data_exchange_async(int32_t SrcDeviceId, void *SrcPtr,
int32_t GenericPluginTy::launch_kernel(int32_t DeviceId, void *TgtEntryPtr,
void **TgtArgs, ptrdiff_t *TgtOffsets,
KernelArgsTy *KernelArgs,
+ KernelExtraArgsTy *KernelExtraArgs,
__tgt_async_info *AsyncInfoPtr) {
auto Err = getDevice(DeviceId).launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets,
- *KernelArgs, AsyncInfoPtr);
+ *KernelArgs, KernelExtraArgs,
+ AsyncInfoPtr);
if (Err) {
REPORT() << "Failure to run target region " << TgtEntryPtr << " in device "
<< DeviceId << ": " << toString(std::move(Err));
diff --git a/offload/plugins-nextgen/common/src/RecordReplay.cpp b/offload/plugins-nextgen/common/src/RecordReplay.cpp
index fc425b16cf9a7..d4ae33c040344 100644
--- a/offload/plugins-nextgen/common/src/RecordReplay.cpp
+++ b/offload/plugins-nextgen/common/src/RecordReplay.cpp
@@ -31,9 +31,10 @@ using namespace error;
RecordReplayTy::InstanceTy::InstanceTy(const GenericKernelTy &Kernel,
uint32_t NumTeams, uint32_t NumThreads,
- uint32_t SharedMemorySize)
+ uint32_t SharedMemorySize,
+ KernelReplayOutcomeTy *ReplayOutcome)
: Kernel(Kernel), NumTeams(NumTeams), NumThreads(NumThreads),
- SharedMemorySize(SharedMemorySize) {
+ SharedMemorySize(SharedMemorySize), ReplayOutcome(ReplayOutcome) {
KernelHash = stable_hash_name(Kernel.getName());
LaunchConfigHash =
stable_hash_combine((stable_hash)NumTeams, (stable_hash)NumThreads,
@@ -92,10 +93,12 @@ Error RecordReplayTy::emitInstanceReport() {
llvm::outs() << "directory: "
<< std::filesystem::absolute(OutputDirectory).string() << "\n";
llvm::outs() << "kernels: " << Instances.size() << "\n";
- for (const auto &Inst : Instances) {
- llvm::outs() << getFilename(Inst, "json", /*IncludeDir=*/false) << ": "
- << Inst.Kernel.getName() << "\n";
- }
+
+ SmallString<128> Filename;
+ for (const auto &Inst : Instances)
+ llvm::outs()
+ << getFilename(Inst, FileTy::Descriptor, /*IncludeDir=*/false).c_str()
+ << ": " << Inst.Kernel.getName() << "\n";
llvm::outs() << "=== record report end ===\n";
return Plugin::success();
}
@@ -103,10 +106,11 @@ Error RecordReplayTy::emitInstanceReport() {
std::pair<const RecordReplayTy::InstanceTy &, bool>
RecordReplayTy::registerInstance(const GenericKernelTy &Kernel,
uint32_t NumTeams, uint32_t NumThreads,
- uint32_t SharedMemorySize) {
+ uint32_t SharedMemorySize,
+ KernelReplayOutcomeTy *ReplayOutcome) {
std::lock_guard<std::mutex> LG(InstancesLock);
- auto [It, Inserted] =
- Instances.emplace(Kernel, NumTeams, NumThreads, SharedMemorySize);
+ auto [It, Inserted] = Instances.emplace(Kernel, NumTeams, NumThreads,
+ SharedMemorySize, ReplayOutcome);
// Increase the number of occurrences.
It->Occurrences += 1;
return {*It, Inserted};
@@ -131,14 +135,16 @@ Error RecordReplayTy::deallocate(void *Ptr) { return Plugin::success(); }
Expected<RecordReplayTy::HandleTy> RecordReplayTy::recordPrologue(
const GenericKernelTy &Kernel, const KernelArgsTy &KernelArgs,
+ const KernelExtraArgsTy *KernelExtraArgs,
const KernelLaunchParamsTy &LaunchParams, uint32_t NumTeams[3],
uint32_t NumThreads[3], uint32_t SharedMemorySize) {
if (!isRecordingOrReplaying())
return HandleTy{nullptr, false};
// Register the instance and avoid recording if it is inactive or replaying.
- auto [Instance, First] =
- registerInstance(Kernel, NumTeams[0], NumThreads[0], SharedMemorySize);
+ auto [Instance, First] = registerInstance(
+ Kernel, NumTeams[0], NumThreads[0], SharedMemorySize,
+ (KernelExtraArgs) ? KernelExtraArgs->ReplayOutcome : nullptr);
HandleTy Handle{&Instance, First};
if (isReplaying() || !First)
@@ -158,28 +164,39 @@ Error RecordReplayTy::recordEpilogue(const GenericKernelTy &Kernel,
if (!shouldRecordEpilogue() || !Handle.Active)
return Plugin::success();
- return recordEpilogueImpl(Kernel, *Handle.Instance);
+ const InstanceTy &Instance = *Handle.Instance;
+ if (auto Err = recordEpilogueImpl(Kernel, Instance))
+ return Err;
+
+ // If necessary, inform the replaying tool about where the epilogue snapshot
+ // file has been stored.
+ if (isReplaying() && Instance.ReplayOutcome) {
+ SmallString<128> Filename = getFilename(Instance, FileTy::EpilogueSnapshot);
+ Instance.ReplayOutcome->OutputFilepath = Filename;
+ }
+ return Plugin::success();
}
Error NativeRecordReplayTy::recordPrologueImpl(
const GenericKernelTy &Kernel, const InstanceTy &Instance,
const KernelArgsTy &KernelArgs, const KernelLaunchParamsTy &LaunchParams) {
- std::string SnapshotFilename = getFilename(Instance, "record_input");
- if (auto Err = recordSnapshot(SnapshotFilename))
+ SmallString<128> SnapshotFilename =
+ getFilename(Instance, FileTy::PrologueSnapshot);
+ if (auto Err = recordSnapshot(SnapshotFilename.c_str()))
return Err;
- std::string GlobalsFilename = getFilename(Instance, "globals");
- if (auto Err = recordGlobals(GlobalsFilename))
+ SmallString<128> GlobalsFilename = getFilename(Instance, FileTy::Globals);
+ if (auto Err = recordGlobals(GlobalsFilename.c_str()))
return Err;
- std::string ImageFilename = getFilename(Instance, "image");
- return recordImage(Kernel, ImageFilename);
+ SmallString<128> ImageFilename = getFilename(Instance, FileTy::Program);
+ return recordImage(Kernel, ImageFilename.c_str());
}
Error NativeRecordReplayTy::recordEpilogueImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance) {
- std::string SnapshotFilename =
- getFilename(Instance, isRecording() ? "record_output" : "replay_output");
+ SmallString<128> SnapshotFilename =
+ getFilename(Instance, FileTy::EpilogueSnapshot);
return recordSnapshot(SnapshotFilename);
}
@@ -207,9 +224,10 @@ Error NativeRecordReplayTy::recordDescImpl(
JsonArgOffsets.push_back(0);
JsonKernelInfo["ArgOffsets"] = json::Value(std::move(JsonArgOffsets));
- std::string JsonFilename = getFilename(Instance, "json");
+ SmallString<128> JsonFilename = getFilename(Instance, FileTy::Descriptor);
+
std::error_code EC;
- raw_fd_ostream JsonOS(JsonFilename, EC);
+ raw_fd_ostream JsonOS(JsonFilename.c_str(), EC);
if (EC)
return Plugin::error(ErrorCode::HOST_IO, "saving kernel json file");
JsonOS << json::Value(std::move(JsonKernelInfo));
@@ -217,17 +235,34 @@ Error NativeRecordReplayTy::recordDescImpl(
return Plugin::success();
}
-std::string NativeRecordReplayTy::getFilenameImpl(const InstanceTy &Instance,
- StringRef Suffix,
- bool IncludeDirectory) {
+StringRef NativeRecordReplayTy::getExtension(FileTy FileType) {
+ switch (FileType) {
+ case FileTy::PrologueSnapshot:
+ return "record_input";
+ case FileTy::EpilogueSnapshot:
+ return isRecording() ? "record_output" : "replay_output";
+ case FileTy::Descriptor:
+ return "json";
+ case FileTy::Globals:
+ return "globals";
+ case FileTy::Program:
+ return "image";
+ }
+ return "";
+}
+
+SmallString<128>
+NativeRecordReplayTy::getFilenameImpl(const InstanceTy &Instance,
+ FileTy FileType, bool IncludeDirectory) {
std::filesystem::path Filepath = IncludeDirectory ? OutputDirectory : "";
Filepath /= std::to_string(Instance.KernelHash) + "_" +
std::to_string(Instance.LaunchConfigHash);
- Filepath.replace_extension(Suffix.data());
- return Filepath.string();
+ Filepath.replace_extension(getExtension(FileType).data());
+ SmallString<128> Filename(Filepath.c_str());
+ return Filename;
}
-Error NativeRecordReplayTy::recordSnapshot(const std::string &Filename) {
+Error NativeRecordReplayTy::recordSnapshot(StringRef Filename) {
// Another thread may be allocating memory. The size can only increase.
AllocationLock.lock();
uint64_t RecordSize = CurrentSize;
@@ -255,7 +290,7 @@ Error NativeRecordReplayTy::recordSnapshot(const std::string &Filename) {
}
Error NativeRecordReplayTy::recordImage(const GenericKernelTy &Kernel,
- const std::string &Filename) {
+ StringRef Filename) {
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
if (EC)
@@ -265,7 +300,7 @@ Error NativeRecordReplayTy::recordImage(const GenericKernelTy &Kernel,
return Plugin::success();
}
-Error NativeRecordReplayTy::recordGlobals(const std::string &Filename) {
+Error NativeRecordReplayTy::recordGlobals(StringRef Filename) {
AllocationLock.lock();
// Copy the globals into a local vector so we can read it safely from this
// thread. This vector should have a few entries in general. No need to lock
diff --git a/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp b/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp
new file mode 100644
index 0000000000000..803b633315527
--- /dev/null
+++ b/offload/test/tools/omp-kernel-replay/record-replay-diff-teams-threads.cpp
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic
+// RUN: rm -rf %t.testdir
+// RUN: mkdir -p %t.testdir
+// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=1 --num-threads=1 {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=2 --num-threads=32 {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-teams=32 --num-threads=64 {}
+// clang-format on
+
+// REQUIRES: gpu
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: intelgpu
+
+#include <cstdint>
+#include <cstdio>
+
+int main() {
+ size_t Size = 1000;
+ uint64_t *Data = new uint64_t[Size];
+
+ for (size_t I = 0; I < Size; ++I) {
+ Data[I] = 20;
+ }
+
+#pragma omp target teams distribute parallel for thread_limit(128) \
+ map(tofrom : Data[0 : Size])
+ for (size_t I = 0; I < Size; ++I) {
+ Data[I] = 10 + (uint64_t)I;
+ }
+
+ uint64_t Sum = 0;
+ for (size_t I = 0; I < Size; ++I) {
+ Sum += Data[I];
+ }
+
+ // CHECK: PASS
+ if (Sum == 509500)
+ printf("PASS\n");
+
+ delete[] Data;
+}
diff --git a/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp b/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp
new file mode 100644
index 0000000000000..9b65c38f98390
--- /dev/null
+++ b/offload/test/tools/omp-kernel-replay/record-replay-diff-threads.cpp
@@ -0,0 +1,46 @@
+// clang-format off
+// RUN: %libomptarget-compilexx-generic
+// RUN: rm -rf %t.testdir
+// RUN: mkdir -p %t.testdir
+// RUN: env LIBOMPTARGET_RECORD=1 LIBOMPTARGET_RECORD_MEMSIZE=536870912 LIBOMPTARGET_RECORD_DIR=%t.testdir %libomptarget-run-generic 2>&1 | %fcheck-generic
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=1 {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=32 {}
+// RUN: ls -t %t.testdir/*.json | sed -n '1p' | grep . | xargs -I {} %omp-kernel-replay --verify --num-threads=64 {}
+// clang-format on
+
+// REQUIRES: gpu
+
+// UNSUPPORTED: aarch64-unknown-linux-gnu
+// UNSUPPORTED: x86_64-unknown-linux-gnu
+// UNSUPPORTED: s390x-ibm-linux-gnu
+// UNSUPPORTED: intelgpu
+
+#include <cstdint>
+#include <cstdio>
+
+int main() {
+ size_t Size = 1000;
+ uint64_t *Data = new uint64_t[Size];
+
+ for (size_t I = 0; I < Size; ++I) {
+ Data[I] = 20;
+ }
+
+#pragma omp target teams distribute parallel for num_teams(256) \
+ thread_limit(128) map(tofrom : Data[0 : Size])
+ for (size_t I = 0; I < Size; ++I) {
+ Data[I] = 10 + (uint64_t)I;
+ }
+
+ uint64_t Sum = 0;
+ for (size_t I = 0; I < Size; ++I) {
+ Sum += Data[I];
+ }
+
+ // CHECK: PASS
+ if (Sum == 509500)
+ printf("PASS\n");
+
+ delete[] Data;
+}
diff --git a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
index 1d364bbe2418f..3f22c7c223fa4 100644
--- a/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
+++ b/offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
@@ -197,17 +197,22 @@ int main(int argc, char **argv) {
const_cast<char *>(DeviceMemoryMB.get()->getBuffer().data()),
DeviceMemoryMB.get()->getBufferSize());
+ KernelReplayOutcomeTy ReplayOutcome;
+
Rc = __tgt_target_kernel_replay(
/*Loc=*/nullptr, DeviceId, OffloadEntries[0].Address,
(char *)RecordedData, DeviceMemoryMB.get()->getBufferSize(),
NumGlobals ? &OffloadEntries[1] : nullptr, NumGlobals, TgtArgs.data(),
TgtArgOffsets.data(), NumArgs.value(), NumTeams, NumThreads,
- SharedMemorySize, LoopTripCount.value());
+ SharedMemorySize, LoopTripCount.value(), &ReplayOutcome);
if (Rc != OMP_TGT_SUCCESS)
reportFatalUsageError("Error replaying kernel");
int ErrorDetected = 0;
if (VerifyOpt) {
+ if (ReplayOutcome.OutputFilepath.empty())
+ reportFatalUsageError("Replay output file was not generated");
+
Filepath.replace_extension("record_output");
ErrorOr<std::unique_ptr<MemoryBuffer>> OriginalOutputMB =
MemoryBuffer::getFile(Filepath.string(),
@@ -218,9 +223,8 @@ int main(int argc, char **argv) {
"Error reading the kernel record output file. Make sure "
"LIBOMPTARGET_RECORD_OUTPUT is set when recording");
- Filepath.replace_extension("replay_output");
ErrorOr<std::unique_ptr<MemoryBuffer>> ReplayOutputMB =
- MemoryBuffer::getFile(Filepath.string(),
+ MemoryBuffer::getFile(ReplayOutcome.OutputFilepath.c_str(),
/*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!ReplayOutputMB)
More information about the llvm-commits
mailing list