Cold-Air-Networks/mirror-llvm-project
2
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions 7

[offload] Fix kernel record/replay and add extensible mechanism (#190588)

Browse source 
This commit fixes the kernel record replay on both AMD and CUDA devices. It
also re-organizes the record replay code, moves the whole code to separate
files, and makes it extensible to support other record formats (potentially in
the future). The environment variables for controlling the recording have also
been modified.
This commit is contained in:
Kevin Sala Penades 2026-04-15 18:23:10 -07:00 committed by GitHub
parent 4a24c685fe
commit 12a4adf6f2
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
21 changed files with 1282 additions and 642 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

1
offload/cmake/OpenMPTesting.cmake
View file

@ -12,6 +12,7 @@ endif()
set(OFFLOAD_NOT_EXECUTABLE ${LLVM_TOOLS_BINARY_DIR}/not)
set(OFFLOAD_DEVICE_INFO_EXECUTABLE ${LLVM_RUNTIME_OUTPUT_INTDIR}/llvm-offload-device-info)
set(OFFLOAD_TBLGEN_EXECUTABLE ${LLVM_RUNTIME_OUTPUT_INTDIR}/offload-tblgen)
set(OMP_KERNEL_REPLAY ${LLVM_RUNTIME_OUTPUT_INTDIR}/llvm-omp-kernel-replay)
# Set the information that we know.
set(OPENMP_TEST_COMPILER_ID "Clang")

11
offload/include/omptarget.h
View file

@ -428,9 +428,12 @@ void __tgt_target_nowait_query(void **AsyncHandle);
/// device memory.
int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId, void *HostPtr,
void *DeviceMemory, int64_t DeviceMemorySize,
void **TgtArgs, ptrdiff_t *TgtOffsets,
int32_t NumArgs, int32_t NumTeams,
int32_t ThreadLimit, uint64_t LoopTripCount);
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);
void __tgt_set_info_flag(uint32_t);
@ -438,7 +441,7 @@ int __tgt_print_device_info(int64_t DeviceId);
int __tgt_activate_record_replay(int64_t DeviceId, uint64_t MemorySize,
void *VAddr, bool IsRecord, bool SaveOutput,
uint64_t &ReqPtrArgOffset);
bool EmitReport, const char *OutputDirPath);
// Registers a callback for the RPC server. Expects this function type.
// unsigned callback(rpc::Server::Port *Port, unsigned NumLanes). See the RPC

21
offload/libomptarget/device.cpp
View file

@ -89,12 +89,23 @@ llvm::Error DeviceTy::init() {
// Enables recording kernels if set.
BoolEnvar OMPX_RecordKernel("LIBOMPTARGET_RECORD", false);
if (OMPX_RecordKernel) {
// Enables saving the device memory kernel output post execution if set.
BoolEnvar OMPX_ReplaySaveOutput("LIBOMPTARGET_RR_SAVE_OUTPUT", false);
BoolEnvar OMPX_RecordOutput("LIBOMPTARGET_RECORD_OUTPUT", true);
Int64Envar OMPX_RecordMemSize("LIBOMPTARGET_RECORD_MEMSIZE",
8 * 1024 * 1024 * 1024ULL);
Int32Envar OMPX_RecordDevice("LIBOMPTARGET_RECORD_DEVICE", 0);
StringEnvar OMPX_RecordOutputDir("LIBOMPTARGET_RECORD_DIR", "");
BoolEnvar OMPX_EmitRecordReport("LIBOMPTARGET_RECORD_REPORT", false);
if (OMPX_RecordDevice != RTLDeviceID)
return llvm::Error::success();
uint64_t ReqPtrArgOffset;
RTL->initialize_record_replay(RTLDeviceID, 0, nullptr, true,
OMPX_ReplaySaveOutput, ReqPtrArgOffset);
Ret = RTL->initialize_record_replay(
RTLDeviceID, OMPX_RecordMemSize, nullptr,
/*IsRecord=*/true, /*IsNative=*/true, OMPX_RecordOutput,
OMPX_EmitRecordReport, OMPX_RecordOutputDir.get().c_str());
if (Ret != OFFLOAD_SUCCESS)
return error::createOffloadError(error::ErrorCode::BACKEND_FAILURE,
"failed to initialize RR in device %d\n",
DeviceID);
}
return llvm::Error::success();

53
offload/libomptarget/interface.cpp
View file

@ -473,25 +473,32 @@ EXTERN int __tgt_target_kernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
/// Activates the record replay mechanism.
/// \param DeviceId The device identifier to execute the target region.
/// \param MemorySize The number of bytes to be (pre-)allocated
/// by the bump allocator
/// by the record replay allocator.
/// /param IsRecord Activates the record replay mechanism in
/// 'record' mode or 'replay' mode.
/// 'record' or 'replay' mode.
/// /param SaveOutput Store the device memory after kernel
/// execution on persistent storage
/// execution on persistent storage.
/// /param EmitReport Emit a summary report after the recording.
/// /param OutputDirPath The output directory where the record replay files
/// should be stored. An empty string or nullptr indicates the current working
/// directory should be used.
EXTERN int __tgt_activate_record_replay(int64_t DeviceId, uint64_t MemorySize,
void *VAddr, bool IsRecord,
bool SaveOutput,
uint64_t &ReqPtrArgOffset) {
bool SaveOutput, bool EmitReport,
const char *OutputDirPath) {
assert(PM && "Runtime not initialized");
OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
auto DeviceOrErr = PM->getDevice(DeviceId);
if (!DeviceOrErr)
FATAL_MESSAGE(DeviceId, "%s", toString(DeviceOrErr.takeError()).c_str());
[[maybe_unused]] int Rc = target_activate_rr(
*DeviceOrErr, MemorySize, VAddr, IsRecord, SaveOutput, ReqPtrArgOffset);
assert(Rc == OFFLOAD_SUCCESS &&
"__tgt_activate_record_replay unexpected failure!");
int Rc = target_activate_rr(*DeviceOrErr, MemorySize, VAddr, IsRecord,
SaveOutput, EmitReport, OutputDirPath);
if (Rc != OFFLOAD_SUCCESS) {
ODBG(ODT_Interface) << "Record replay failed to activate in device "
<< DeviceId;
return OMP_TGT_FAIL;
}
return OMP_TGT_SUCCESS;
}
@ -512,12 +519,12 @@ EXTERN int __tgt_activate_record_replay(int64_t DeviceId, uint64_t MemorySize,
/// execution.
/// \param LoopTripCount The pre-recorded value of the loop tripcount, if any.
/// \return OMP_TGT_SUCCESS on success, OMP_TGT_FAIL on failure.
EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
void *HostPtr, void *DeviceMemory,
int64_t DeviceMemorySize, void **TgtArgs,
ptrdiff_t *TgtOffsets, int32_t NumArgs,
int32_t NumTeams, int32_t ThreadLimit,
uint64_t LoopTripCount) {
EXTERN 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) {
assert(PM && "Runtime not initialized");
OMPT_IF_BUILT(ReturnAddressSetterRAII RA(__builtin_return_address(0)));
if (checkDevice(DeviceId, Loc)) {
@ -534,14 +541,18 @@ EXTERN int __tgt_target_kernel_replay(ident_t *Loc, int64_t DeviceId,
/*CodePtr=*/OMPT_GET_RETURN_ADDRESS);)
AsyncInfoTy AsyncInfo(*DeviceOrErr);
int Rc = target_replay(Loc, *DeviceOrErr, HostPtr, DeviceMemory,
DeviceMemorySize, TgtArgs, TgtOffsets, NumArgs,
NumTeams, ThreadLimit, LoopTripCount, AsyncInfo);
int Rc =
target_replay(Loc, *DeviceOrErr, HostPtr, DeviceMemory, DeviceMemorySize,
Globals, NumGlobals, TgtArgs, TgtOffsets, NumArgs, NumTeams,
ThreadLimit, SharedMemorySize, LoopTripCount, AsyncInfo);
if (Rc == OFFLOAD_SUCCESS)
Rc = AsyncInfo.synchronize();
handleTargetOutcome(Rc == OFFLOAD_SUCCESS, Loc);
assert(Rc == OFFLOAD_SUCCESS &&
"__tgt_target_kernel_replay unexpected failure!");
if (Rc != OFFLOAD_SUCCESS) {
ODBG(ODT_Interface) << "Kernel replay failed in device " << DeviceId;
return OMP_TGT_FAIL;
}
return OMP_TGT_SUCCESS;
}

102
offload/libomptarget/omptarget.cpp
View file

@ -2379,67 +2379,105 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
/// and informing the record-replayer of whether to store the output
/// in some file.
int target_activate_rr(DeviceTy &Device, uint64_t MemorySize, void *VAddr,
bool IsRecord, bool SaveOutput,
uint64_t &ReqPtrArgOffset) {
return Device.RTL->initialize_record_replay(Device.DeviceID, MemorySize,
VAddr, IsRecord, SaveOutput,
ReqPtrArgOffset);
bool IsRecord, bool SaveOutput, bool EmitReport,
const char *OutputDirPath) {
return Device.RTL->initialize_record_replay(
Device.DeviceID, MemorySize, VAddr, IsRecord,
/*IsNative=*/true, SaveOutput, EmitReport, OutputDirPath);
}
/// Executes a kernel using pre-recorded information for loading to
/// device memory to launch the target kernel with the pre-recorded
/// configuration.
int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
void *DeviceMemory, int64_t DeviceMemorySize, void **TgtArgs,
ptrdiff_t *TgtOffsets, int32_t NumArgs, int32_t NumTeams,
int32_t ThreadLimit, uint64_t LoopTripCount,
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,
AsyncInfoTy &AsyncInfo) {
int32_t DeviceId = Device.DeviceID;
TableMap *TM = getTableMap(HostPtr);
// Fail if the table map fails to find the target kernel pointer for the
// provided host pointer.
if (!TM) {
REPORT() << "Host ptr " << HostPtr
<< " does not have a matching target pointer.";
return OFFLOAD_FAIL;
int32_t NumSymbols = NumGlobals + 1;
struct SymbolDataTy {
void *DevPtr = nullptr;
TableMap *TM = nullptr;
__tgt_target_table *TargetTable = nullptr;
};
SmallVector<SymbolDataTy> Symbols(NumSymbols);
for (int32_t I = 0; I < NumSymbols; ++I) {
// The first symbol is the kernel entry.
void *SymbolHostPtr = (I == 0) ? HostPtr : Globals[I - 1].Address;
// Get the table map for each symbol.
Symbols[I].TM = getTableMap(SymbolHostPtr);
if (!Symbols[I].TM) {
REPORT() << "Host pointer " << SymbolHostPtr
<< " does not have a matching target pointer.";
return OFFLOAD_FAIL;
}
}
// Retrieve the target table of offloading entries.
__tgt_target_table *TargetTable = nullptr;
// Retrieve the target table for each symbol.
{
std::lock_guard<std::mutex> TrlTblLock(PM->TrlTblMtx);
assert(TM->Table->TargetsTable.size() > (size_t)DeviceId &&
"Not expecting a device ID outside the table's bounds!");
TargetTable = TM->Table->TargetsTable[DeviceId];
for (auto &S : Symbols) {
S.TargetTable = S.TM->Table->TargetsTable[DeviceId];
assert(S.TargetTable && "Global data has not been mapped\n");
}
}
assert(TargetTable && "Global data has not been mapped\n");
// Retrieve the target kernel pointer, allocate and store the recorded device
// memory data, and launch device execution.
void *TgtEntryPtr = TargetTable->EntriesBegin[TM->Index].Address;
ODBG(ODT_Kernel) << "Launching target execution "
<< TargetTable->EntriesBegin[TM->Index].SymbolName
<< " with pointer " << TgtEntryPtr << " (index=" << TM->Index
<< ").";
// Retrieve the device pointers for each symbol.
for (auto &S : Symbols)
S.DevPtr = S.TargetTable->EntriesBegin[S.TM->Index].Address;
// Initialize the device memory of each global.
for (int32_t I = 0; I < NumGlobals; ++I) {
assert(Globals[I].AuxAddr && "Global has no AuxAddr.");
// Initialize the value of the global in the device.
int Ret = Device.submitData(Symbols[I + 1].DevPtr, Globals[I].AuxAddr,
Globals[I].Size, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Failed to submit data to a global.";
return OFFLOAD_FAIL;
}
}
void *TgtPtr = Device.allocData(DeviceMemorySize, /*HstPtr=*/nullptr,
TARGET_ALLOC_DEFAULT);
Device.submitData(TgtPtr, DeviceMemory, DeviceMemorySize, AsyncInfo);
if (!TgtPtr) {
REPORT() << "Failed to allocate device memory.";
return OFFLOAD_FAIL;
}
int Ret =
Device.submitData(TgtPtr, DeviceMemory, DeviceMemorySize, AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Failed to submit data to a global.";
return OFFLOAD_FAIL;
}
KernelArgsTy KernelArgs{};
KernelArgs.Version = OMP_KERNEL_ARG_VERSION;
KernelArgs.NumArgs = NumArgs;
KernelArgs.Tripcount = LoopTripCount;
KernelArgs.NumTeams[0] = NumTeams;
KernelArgs.NumTeams[1] = 1;
KernelArgs.NumTeams[2] = 1;
KernelArgs.ThreadLimit[0] = ThreadLimit;
KernelArgs.ThreadLimit[1] = 1;
KernelArgs.ThreadLimit[2] = 1;
KernelArgs.DynCGroupMem = SharedMemorySize;
int Ret = Device.launchKernel(TgtEntryPtr, TgtArgs, TgtOffsets, KernelArgs,
AsyncInfo);
Ret = Device.launchKernel(Symbols[0].DevPtr, TgtArgs, TgtOffsets, KernelArgs,
AsyncInfo);
if (Ret != OFFLOAD_SUCCESS) {
REPORT() << "Executing target region abort target.";
REPORT() << "Failed to launch kernel replay.";
return OFFLOAD_FAIL;
}
return OFFLOAD_SUCCESS;
}

11
offload/libomptarget/private.h
View file

@ -27,14 +27,17 @@ extern int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
KernelArgsTy &KernelArgs, AsyncInfoTy &AsyncInfo);
extern int target_activate_rr(DeviceTy &Device, uint64_t MemorySize,
void *ReqAddr, bool isRecord, bool SaveOutput,
uint64_t &ReqPtrArgOffset);
void *ReqAddr, bool IsRecord, bool SaveOutput,
bool EmitReport, const char *OutputDirPath);
extern int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
void *DeviceMemory, int64_t DeviceMemorySize,
void **TgtArgs, ptrdiff_t *TgtOffsets, int32_t NumArgs,
const llvm::offloading::EntryTy *Globals,
int32_t NumGlobals, void **TgtArgs,
ptrdiff_t *TgtOffsets, int32_t NumArgs,
int32_t NumTeams, int32_t ThreadLimit,
uint64_t LoopTripCount, AsyncInfoTy &AsyncInfo);
uint32_t SharedMemorySize, uint64_t LoopTripCount,
AsyncInfoTy &AsyncInfo);
extern void handleTargetOutcome(bool Success, ident_t *Loc);

7
offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.cpp
View file

@ -75,6 +75,13 @@ DLWRAP(hsa_amd_profiling_set_profiler_enabled, 2)
DLWRAP(hsa_code_object_reader_create_from_memory, 3)
DLWRAP(hsa_code_object_reader_destroy, 1)
DLWRAP(hsa_executable_load_agent_code_object, 5)
DLWRAP(hsa_amd_vmem_address_reserve, 4)
DLWRAP(hsa_amd_vmem_address_free, 2)
DLWRAP(hsa_amd_vmem_handle_create, 5)
DLWRAP(hsa_amd_vmem_handle_release, 1)
DLWRAP(hsa_amd_vmem_map, 5)
DLWRAP(hsa_amd_vmem_unmap, 2)
DLWRAP(hsa_amd_vmem_set_access, 4)
DLWRAP_FINALIZE()

7
offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa.h
View file

@ -116,6 +116,13 @@ typedef struct hsa_isa_s {
uint64_t handle;
} hsa_isa_t;
typedef enum {
HSA_ACCESS_PERMISSION_NONE = 0,
HSA_ACCESS_PERMISSION_RO = 1,
HSA_ACCESS_PERMISSION_WO = 2,
HSA_ACCESS_PERMISSION_RW = 3
} hsa_access_permission_t;
hsa_status_t hsa_system_get_info(hsa_system_info_t attribute, void *value);
hsa_status_t hsa_agent_get_info(hsa_agent_t agent, hsa_agent_info_t attribute,

37
offload/plugins-nextgen/amdgpu/dynamic_hsa/hsa_ext_amd.h
View file

@ -163,6 +163,20 @@ typedef struct hsa_amd_pointer_info_s {
size_t sizeInBytes;
} hsa_amd_pointer_info_t;
typedef enum {
MEMORY_TYPE_NONE,
MEMORY_TYPE_PINNED,
} hsa_amd_memory_type_t;
typedef struct hsa_amd_vmem_alloc_handle_s {
uint64_t handle;
} hsa_amd_vmem_alloc_handle_t;
typedef struct hsa_amd_memory_access_desc_s {
hsa_access_permission_t permissions;
hsa_agent_t agent_handle;
} hsa_amd_memory_access_desc_t;
hsa_status_t hsa_amd_pointer_info(const void* ptr,
hsa_amd_pointer_info_t* info,
void* (*alloc)(size_t),
@ -181,6 +195,29 @@ hsa_amd_profiling_get_dispatch_time(hsa_agent_t agent, hsa_signal_t signal,
hsa_status_t hsa_amd_profiling_set_profiler_enabled(hsa_queue_t *queue,
int enable);
hsa_status_t hsa_amd_vmem_address_reserve(void **va, size_t size,
uint64_t address, uint64_t flags);
hsa_status_t hsa_amd_vmem_address_free(void *va, size_t size);
hsa_status_t
hsa_amd_vmem_handle_create(hsa_amd_memory_pool_t pool, size_t size,
hsa_amd_memory_type_t type, uint64_t flags,
hsa_amd_vmem_alloc_handle_t *memory_handle);
hsa_status_t
hsa_amd_vmem_handle_release(hsa_amd_vmem_alloc_handle_t memory_handle);
hsa_status_t hsa_amd_vmem_map(void *va, size_t size, size_t in_offset,
hsa_amd_vmem_alloc_handle_t memory_handle,
uint64_t flags);
hsa_status_t hsa_amd_vmem_unmap(void *va, size_t size);
hsa_status_t hsa_amd_vmem_set_access(void *va, size_t size,
const hsa_amd_memory_access_desc_t *desc,
size_t desc_cnt);
#ifdef __cplusplus
}
#endif

101
offload/plugins-nextgen/amdgpu/src/rtl.cpp
View file

@ -290,7 +290,7 @@ struct AMDGPUMemoryPoolTy {
if (auto Err = getAttr(HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, GlobalFlags))
return Err;
return Plugin::success();
return getAttr(HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, Granule);
}
/// Getter of the HSA memory pool.
@ -320,6 +320,9 @@ struct AMDGPUMemoryPoolTy {
return (GlobalFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_KERNARG_INIT);
}
/// Get the allocation granularity of the pool.
size_t getGranule() const { return Granule; }
/// Allocate memory on the memory pool.
Error allocate(size_t Size, void **PtrStorage) {
hsa_status_t Status =
@ -400,6 +403,9 @@ private:
/// The global flags of memory pool. Only valid if the memory pool belongs to
/// the global segment.
uint32_t GlobalFlags;
/// The page size in this memory pool.
size_t Granule;
};
/// Class that implements a memory manager that gets memory from a specific
@ -2320,6 +2326,91 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
return Plugin::success();
}
/// Suggest a virtual address for device memory mapping.
void *getSuggestedVirtualAddress() override {
return reinterpret_cast<void *>(0x1534f7e00000ULL);
}
/// Allocate \p Size bytes on the device and hints the backend to map it to
/// virtual address \p VAddr. The function returns the allocated virtual
/// address. The memory must be deallocated through
/// GenericDeviceTy::deallocateWithVirtualAddress().
Expected<void *> allocateWithVirtualAddress(uint64_t Size,
void *VAddr) override {
uint64_t ExpectedVAddr = 0;
if (VAddr != nullptr)
ExpectedVAddr = static_cast<uint64_t>(reinterpret_cast<uintptr_t>(VAddr));
// Transparently round up to a multiple of the page size.
auto *Pool = CoarseGrainedMemoryPools[0];
Size = utils::roundUp(Size, (uint64_t)Pool->getGranule());
// Reserve the virtual address range.
hsa_status_t Status =
hsa_amd_vmem_address_reserve(&VAddr, Size, ExpectedVAddr, 0);
if (auto Err = Plugin::check(Status,
"error in hsa_amd_vmem_address_reserve: %s\n"))
return Err;
if (ExpectedVAddr != 0 && reinterpret_cast<void *>(ExpectedVAddr) != VAddr)
ODBG(OLDT_Alloc)
<< "hsa_amd_vmem_address_reserve reserved device virtual address "
<< VAddr << " instead of " << reinterpret_cast<void *>(ExpectedVAddr);
// Create a handle of the allocation.
hsa_amd_vmem_alloc_handle_t Handle;
Status = hsa_amd_vmem_handle_create(Pool->get(), Size, MEMORY_TYPE_PINNED,
0, &Handle);
if (auto Err =
Plugin::check(Status, "error in hsa_amd_vmem_handle_create: %s\n"))
return Err;
// Map the virtual address range to the memory allocation.
Status = hsa_amd_vmem_map(VAddr, Size, 0, Handle, 0);
if (auto Err = Plugin::check(Status, "error in hsa_amd_vmem_map: %s\n"))
return Err;
// Set the memory access properties for the allocation.
hsa_amd_memory_access_desc_t Desc;
Desc.agent_handle = Agent;
Desc.permissions = HSA_ACCESS_PERMISSION_RW;
Status = hsa_amd_vmem_set_access(VAddr, Size, &Desc, 1);
if (auto Err =
Plugin::check(Status, "error in hsa_amd_vmem_set_access: %s\n"))
return Err;
// Register the virtual address range in the tracker.
if (auto Err = VMemTracker.registerReservation(VAddr, Size, Handle))
return Err;
return VAddr;
}
/// Deallocate device memory \p VAddr, which was allocated through
/// GenericDeviceTy::allocateWithVirtualAddress(), and unmap the virtual
/// address range.
Error deallocateWithVirtualAddress(void *VAddr, uint64_t) override {
// Unregister the virtual address range and obtain the information about
// the reservation.
auto InfoOrErr = VMemTracker.unregisterReservation(VAddr);
if (!InfoOrErr)
return InfoOrErr.takeError();
auto [Size, Handle] = *InfoOrErr;
hsa_status_t Status = hsa_amd_vmem_unmap(VAddr, Size);
if (auto Err = Plugin::check(Status, "error in hsa_amd_vmem_unmap: %s\n"))
return Err;
Status = hsa_amd_vmem_handle_release(Handle);
if (auto Err =
Plugin::check(Status, "error in hsa_amd_vmem_handle_release: %s\n"))
return Err;
Status = hsa_amd_vmem_address_free(VAddr, Size);
return Plugin::check(Status, "error in hsa_amd_vmem_address_free: %s\n");
}
Error unloadBinaryImpl(DeviceImageTy *Image) override {
AMDGPUDeviceImageTy &AMDImage = static_cast<AMDGPUDeviceImageTy &>(*Image);
@ -3204,10 +3295,7 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy {
if (Status == HSA_STATUS_SUCCESS)
PoolNode.add("Allocatable", TmpBool);
Status = Pool->getAttrRaw(HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
TmpSt);
if (Status == HSA_STATUS_SUCCESS)
PoolNode.add("Runtime Alloc Granule", TmpSt, "bytes");
PoolNode.add("Runtime Alloc Granule", Pool->getGranule(), "bytes");
Status = Pool->getAttrRaw(
HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENT, TmpSt);
@ -3502,6 +3590,9 @@ private:
/// True is the system is configured with XNACK-Enabled.
/// False otherwise.
bool IsXnackEnabled = false;
/// Tracker for virtual address reservations.
VMemTrackerTy<hsa_amd_vmem_alloc_handle_t> VMemTracker;
};
Error AMDGPUDeviceImageTy::loadExecutable(const AMDGPUDeviceTy &Device) {

1
offload/plugins-nextgen/common/CMakeLists.txt
View file

@ -13,6 +13,7 @@ add_library(PluginCommon OBJECT
src/PluginInterface.cpp
src/GlobalHandler.cpp
src/JIT.cpp
src/RecordReplay.cpp
src/RPC.cpp
src/OffloadError.cpp
src/Utils/ELF.cpp

111
offload/plugins-nextgen/common/include/PluginInterface.h
View file

@ -34,13 +34,18 @@
#include "MemoryManager.h"
#include "OffloadError.h"
#include "RPC.h"
#include "RecordReplay.h"
#include "omptarget.h"
#ifdef OMPT_SUPPORT
#include "omp-tools.h"
#endif
#include "llvm/ADT/DenseMapInfo.h"
#include "llvm/ADT/DenseSet.h"
#include "llvm/ADT/Hashing.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StableHashing.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/Frontend/OpenMP/OMPGridValues.h"
#include "llvm/Support/Allocator.h"
@ -61,7 +66,6 @@ namespace plugin {
struct GenericPluginTy;
struct GenericKernelTy;
struct GenericDeviceTy;
struct RecordReplayTy;
template <typename ResourceRef> class GenericDeviceResourceManagerTy;
namespace Plugin {
@ -313,6 +317,45 @@ struct DynBlockMemConfTy {
void *FallbackPtr = nullptr;
};
/// Tracker of virtual memory address reservations.
template <typename HandleTy> class VMemTrackerTy {
struct EntryTy {
uint64_t Size;
HandleTy Handle;
};
/// Map of virtual memory address reservations.
DenseMap<void *, EntryTy> VMemMap;
/// Mutex for safe access to the map.
std::mutex Mutex;
public:
/// Register a new virtual address reservation.
Error registerReservation(void *VAddr, uint64_t Size, HandleTy Handle) {
std::lock_guard<std::mutex> Lock(Mutex);
auto It = VMemMap.find(VAddr);
if (It != VMemMap.end())
return Plugin::error(error::ErrorCode::INVALID_ARGUMENT,
"virtual address already reserved");
VMemMap[VAddr] = {Size, Handle};
return Plugin::success();
}
/// Unregister a virtual address reservation and return its information.
Expected<std::pair<uint64_t, HandleTy>> unregisterReservation(void *VAddr) {
std::lock_guard<std::mutex> Lock(Mutex);
auto It = VMemMap.find(VAddr);
if (It == VMemMap.end())
return Plugin::error(error::ErrorCode::INVALID_ARGUMENT,
"virtual address not reserved");
uint64_t Size = It->second.Size;
HandleTy Handle = It->second.Handle;
VMemMap.erase(It);
return std::make_pair(Size, Handle);
}
};
/// Class wrapping a __tgt_device_image and its offload entry table on a
/// specific device. This class is responsible for storing and managing
/// the offload entries for an image on a device.
@ -372,7 +415,8 @@ struct GenericKernelTy {
/// one used to initialize the kernel.
Error launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
AsyncInfoWrapperTy &AsyncInfoWrapper) const;
AsyncInfoWrapperTy &AsyncInfoWrapper,
RecordReplayTy::HandleTy *RRHandle = nullptr) const;
virtual Error launchImpl(GenericDeviceTy &GenericDevice,
uint32_t NumThreads[3], uint32_t NumBlocks[3],
uint32_t DynBlockMemSize, KernelArgsTy &KernelArgs,
@ -778,6 +822,27 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
GenericDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId, int32_t NumDevices,
const llvm::omp::GV &GridValues);
/// Suggest a virtual address for device memory mapping.
virtual void *getSuggestedVirtualAddress() { return nullptr; }
/// Allocate \p Size bytes on the device and hints the backend to map it to
/// virtual address \p VAddr. The function returns the allocated virtual
/// address. The memory must be deallocated through
/// GenericDeviceTy::deallocateWithVirtualAddress().
virtual Expected<void *> allocateWithVirtualAddress(uint64_t Size,
void *VAddr = nullptr) {
return Plugin::error(error::ErrorCode::UNSUPPORTED,
"allocate with virtual address not supported");
}
/// Deallocate device memory \p VAddr, which was allocated through
/// GenericDeviceTy::allocateWithVirtualAddress(), and unmap the virtual
/// address range.
virtual Error deallocateWithVirtualAddress(void *VAddr, uint64_t Size) {
return Plugin::error(error::ErrorCode::UNSUPPORTED,
"allocate with virtual address not supported");
}
/// Get the device identifier within the corresponding plugin. Notice that
/// this id is not unique between different plugins; they may overlap.
int32_t getDeviceId() const { return DeviceId; }
@ -1156,6 +1221,29 @@ struct GenericDeviceTy : public DeviceAllocatorTy {
return ATI;
}
Error initRecordReplay(int64_t Size, void *VAddr, bool IsRecord,
bool IsNative, bool SaveOutput, bool EmitReport,
const char *OutputDirPath) {
if (RecordReplay)
return Plugin::error(error::ErrorCode::INVALID_ARGUMENT,
"RR already initialized");
// Other formats could be supported in the future.
if (!IsNative)
return Plugin::error(error::ErrorCode::UNSUPPORTED,
"non-native RR not available");
RecordReplayTy::StatusTy Status = IsRecord
? RecordReplayTy::StatusTy::Recording
: RecordReplayTy::StatusTy::Replaying;
RecordReplay =
new NativeRecordReplayTy(Status, OutputDirPath ? OutputDirPath : "",
SaveOutput, EmitReport, *this);
return RecordReplay->init(Size, VAddr);
}
RecordReplayTy *getRecordReplay() { return RecordReplay; }
/// Map to record kernel have been launchedl, for error reporting purposes.
ProtectedObj<KernelTraceInfoRecordTy> KernelLaunchTraces;
@ -1218,6 +1306,9 @@ private:
/// Indicate whether failures when locking mapped buffers should be ignored.
bool IgnoreLockMappedFailures;
/// Record and replay manager.
RecordReplayTy *RecordReplay = nullptr;
protected:
/// Environment variables defined by the LLVM OpenMP implementation
/// regarding the initial number of streams and events.
@ -1281,8 +1372,7 @@ struct GenericPluginTy {
/// Construct a plugin instance.
GenericPluginTy(Triple::ArchType TA)
: GlobalHandler(nullptr), JIT(TA), RPCServer(nullptr),
RecordReplay(nullptr) {}
: GlobalHandler(nullptr), JIT(TA), RPCServer(nullptr) {}
virtual ~GenericPluginTy() {}
@ -1374,11 +1464,6 @@ struct GenericPluginTy {
virtual Error deinitRPCDoorbell() { return Plugin::success(); }
/// Get a reference to the record and replay interface for the plugin.
RecordReplayTy &getRecordReplay() {
assert(RecordReplay && "RR interface not initialized");
return *RecordReplay;
}
/// Initialize a device within the plugin.
Error initDevice(int32_t DeviceId);
@ -1471,8 +1556,9 @@ public:
/// Initializes the record and replay mechanism inside the plugin.
int32_t initialize_record_replay(int32_t DeviceId, int64_t MemorySize,
void *VAddr, bool isRecord, bool SaveOutput,
uint64_t &ReqPtrArgOffset);
void *VAddr, bool IsRecord, bool IsNative,
bool SaveOutput, bool EmitReport,
const char *OutputDirPath);
/// Loads the associated binary into the plugin and returns a handle to it.
int32_t load_binary(int32_t DeviceId, __tgt_device_image *TgtImage,
@ -1648,9 +1734,6 @@ private:
/// The interface between the plugin and the GPU for host services.
RPCServerTy *RPCServer;
/// The interface between the plugin and the GPU for host services.
RecordReplayTy *RecordReplay;
};
/// Auxiliary interface class for GenericDeviceResourceManagerTy. This class

258
offload/plugins-nextgen/common/include/RecordReplay.h Normal file
View file

@ -0,0 +1,258 @@
//===- RecordReplay.h - Target independent kernel record replay interface -===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
//===----------------------------------------------------------------------===//
#ifndef OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_RECORDREPLAY_H
#define OPENMP_LIBOMPTARGET_PLUGINS_NEXTGEN_COMMON_RECORDREPLAY_H
#include <cstddef>
#include <cstdint>
#include <filesystem>
#include <mutex>
#include <unordered_set>
#include "Shared/APITypes.h"
#include "Shared/EnvironmentVar.h"
#include "Shared/Utils.h"
#include "OffloadError.h"
#include "llvm/ADT/Hashing.h"
#include "llvm/ADT/SmallVector.h"
#include "llvm/ADT/StableHashing.h"
#include "llvm/Support/Error.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/MemoryBufferRef.h"
#include "llvm/Support/raw_ostream.h"
namespace llvm {
namespace omp {
namespace target {
namespace plugin {
struct GenericKernelTy;
struct GenericDeviceTy;
struct RecordReplayTy {
protected:
struct InstanceTy;
public:
/// Describes the state of the record replay mechanism.
enum StatusTy { Deactivated = 0, Recording, Replaying };
/// Describes the format of the recording and replaying.
enum FormatTy { Native = 0 };
struct HandleTy {
const InstanceTy *Instance = nullptr;
bool Active = false;
};
protected:
/// Address and size of record replay memory space.
void *StartAddr = nullptr;
uint64_t TotalSize = 0;
uint64_t CurrentSize = 0;
/// Status of the record or replay.
StatusTy Status;
/// The path where to store all recorded files.
std::filesystem::path OutputDirectory;
/// Whether a memory snapshot should be recorded a kernel execution.
bool SaveOutput;
/// Whether a report should be emitted afther the recording.
bool EmitReport;
/// Reference to the corresponding device.
GenericDeviceTy &Device;
/// The information for a global.
struct GlobalEntryTy {
std::string Name;
uint64_t Size;
void *Addr;
};
/// List of all globals mapped to the device.
SmallVector<GlobalEntryTy> GlobalEntries;
/// Mutex that protects dynamic allocations and globals.
std::mutex AllocationLock;
// An instance of a kernel record replay.
struct InstanceTy {
/// Reference to the kernel.
const GenericKernelTy &Kernel;
/// The launch configuration parameters.
uint32_t NumTeams = 0;
uint32_t NumThreads = 0;
uint32_t SharedMemorySize = 0;
/// The hashes representing the kernel and the launch configuration.
size_t KernelHash = 0;
size_t LaunchConfigHash = 0;
/// 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);
bool operator==(const InstanceTy &Other) const {
return (KernelHash == Other.KernelHash &&
LaunchConfigHash == Other.LaunchConfigHash &&
NumTeams == Other.NumTeams && NumThreads == Other.NumThreads &&
SharedMemorySize == Other.SharedMemorySize);
}
};
struct InstanceHasher {
std::size_t operator()(const InstanceTy &I) const {
llvm::stable_hash H =
llvm::stable_hash_combine(I.KernelHash, I.LaunchConfigHash);
return static_cast<std::size_t>(H);
}
};
/// Tracker of record replay instances.
std::unordered_set<InstanceTy, InstanceHasher> Instances;
std::mutex InstancesLock;
public:
RecordReplayTy(StatusTy Status, StringRef OutputDirectoryStr, bool SaveOutput,
bool EmitReport, GenericDeviceTy &Device)
: Status(Status), SaveOutput(SaveOutput), EmitReport(EmitReport),
Device(Device) {
if (OutputDirectoryStr == "")
OutputDirectory = std::filesystem::current_path();
else
OutputDirectory = OutputDirectoryStr.data();
}
virtual ~RecordReplayTy() = default;
/// Initialize kernel record replay for the corresponding device.
Error init(uint64_t MemSize, void *VAddr);
Error deinit();
/// Emit the information of each registered instance.
Error emitInstanceReport();
bool isRecording() const { return Status == StatusTy::Recording; }
bool isReplaying() const { return Status == StatusTy::Replaying; }
bool isRecordingOrReplaying() const { return isRecording() || isReplaying(); }
bool shouldRecordPrologue() const { return isRecording(); }
bool shouldRecordEpilogue() const {
return isRecordingOrReplaying() && SaveOutput;
}
/// Add information about a global.
void addGlobal(const char *Name, uint64_t Size, void *Addr) {
std::lock_guard<std::mutex> Lock(AllocationLock);
GlobalEntries.emplace_back(GlobalEntryTy{Name, Size, Addr});
}
/// Register a record replay instance, record the prologue data if necessary,
/// and return the instance's handle. The prologue is the phase just before
/// executing the kernel. This phase can include the recording of memory
/// snapshot, the record descriptor and the globals. When replaying, only the
/// instance is registered.
Expected<HandleTy>
recordPrologue(const GenericKernelTy &Kernel, const KernelArgsTy &KernelArgs,
const KernelLaunchParamsTy &LaunchParams, uint32_t NumTeams[3],
uint32_t NumThreads[3], uint32_t SharedMemorySize);
/// Record the epilogue if necessary, which can include the memory snapshot
/// when recording or replaying.
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);
}
/// Allocates device memory from the record replay space.
Expected<void *> allocate(uint64_t Size);
Error deallocate(void *Ptr);
private:
/// Register an instance and return a reference and whether it was registered
/// as a new instance.
std::pair<const InstanceTy &, bool>
registerInstance(const GenericKernelTy &Kernel, uint32_t NumTeams,
uint32_t NumThreads, uint32_t SharedMemorySize);
/// Record the prologue data.
virtual Error
recordPrologueImpl(const GenericKernelTy &Kernel, const InstanceTy &Instance,
const KernelArgsTy &KernelArgs,
const KernelLaunchParamsTy &LaunchParams) = 0;
/// Record the epilogue data.
virtual Error recordEpilogueImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance) = 0;
/// Record the descriptor of the kernel.
virtual Error recordDescImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance,
const KernelArgsTy &KernelArgs,
const KernelLaunchParamsTy &LaunchParams) = 0;
/// Get a string with the filename.
virtual std::string getFilenameImpl(const InstanceTy &Instance,
StringRef Suffix,
bool IncludeDirectory) = 0;
};
/// The native kernel record replay support.
struct NativeRecordReplayTy : public RecordReplayTy {
NativeRecordReplayTy(StatusTy Status, StringRef OutputDirectoryStr,
bool SaveOutput, bool EmitReport,
GenericDeviceTy &Device)
: RecordReplayTy(Status, OutputDirectoryStr, SaveOutput, EmitReport,
Device) {}
private:
Error recordPrologueImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance,
const KernelArgsTy &KernelArgs,
const KernelLaunchParamsTy &LaunchParams) override;
Error recordEpilogueImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance) override;
Error recordDescImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance,
const KernelArgsTy &KernelArgs,
const KernelLaunchParamsTy &LaunchParams) override;
/// Get a string with the filename.
std::string getFilenameImpl(const InstanceTy &Instance, StringRef Suffix,
bool IncludeDirectory) override;
/// Record a memory snapshot to a file.
Error recordSnapshot(const std::string &Filename);
/// Record the globals to a file.
Error recordGlobals(const std::string &Filename);
/// Record the device image to a file.
Error recordImage(const GenericKernelTy &Kernel, const std::string &Filename);
};
} // namespace plugin
} // namespace target
} // namespace omp
} // namespace llvm
#endif // OPENMP_LIBOMPTARGET_PLUGINS_COMMON_RECORDREPLAY_H

427
offload/plugins-nextgen/common/src/PluginInterface.cpp
View file

@ -29,7 +29,6 @@
#include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/Support/Error.h"
#include "llvm/Support/JSON.h"
#include "llvm/Support/MathExtras.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/Signals.h"
@ -45,339 +44,6 @@ using namespace plugin;
using namespace error;
using namespace llvm::offload::debug;
// TODO: Fix any thread safety issues for multi-threaded kernel recording.
namespace llvm::omp::target::plugin {
struct RecordReplayTy {
// Describes the state of the record replay mechanism.
enum RRStatusTy { RRDeactivated = 0, RRRecording, RRReplaying };
private:
// Memory pointers for recording, replaying memory.
void *MemoryStart = nullptr;
void *MemoryPtr = nullptr;
size_t MemorySize = 0;
size_t TotalSize = 0;
GenericDeviceTy *Device = nullptr;
std::mutex AllocationLock;
RRStatusTy Status = RRDeactivated;
bool ReplaySaveOutput = false;
bool UsedVAMap = false;
uintptr_t MemoryOffset = 0;
// A list of all globals mapped to the device.
struct GlobalEntry {
const char *Name;
uint64_t Size;
void *Addr;
};
llvm::SmallVector<GlobalEntry> GlobalEntries{};
Expected<void *> suggestAddress(uint64_t MaxMemoryAllocation) {
// Get a valid pointer address for this system
auto AddrOrErr =
Device->allocate(1024, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);
if (!AddrOrErr)
return AddrOrErr.takeError();
void *Addr = *AddrOrErr;
if (auto Err = Device->free(Addr))
return std::move(Err);
// Align Address to MaxMemoryAllocation
Addr = (void *)utils::alignPtr((Addr), MaxMemoryAllocation);
return Addr;
}
Error preAllocateVAMemory(uint64_t MaxMemoryAllocation, void *VAddr) {
size_t ASize = MaxMemoryAllocation;
if (!VAddr && isRecording()) {
auto VAddrOrErr = suggestAddress(MaxMemoryAllocation);
if (!VAddrOrErr)
return VAddrOrErr.takeError();
VAddr = *VAddrOrErr;
}
ODBG(OLDT_Alloc) << "Request " << MaxMemoryAllocation
<< " bytes allocated at " << VAddr;
if (auto Err = Device->memoryVAMap(&MemoryStart, VAddr, &ASize))
return Err;
if (isReplaying() && VAddr != MemoryStart) {
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"record-Replay cannot assign the"
"requested recorded address (%p, %p)",
VAddr, MemoryStart);
}
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),
"Allocated %" PRIu64 " bytes at %p for replay.\n", ASize, MemoryStart);
MemoryPtr = MemoryStart;
MemorySize = 0;
TotalSize = ASize;
UsedVAMap = true;
return Plugin::success();
}
Error preAllocateHeuristic(uint64_t MaxMemoryAllocation,
uint64_t RequiredMemoryAllocation, void *VAddr) {
const size_t MAX_MEMORY_ALLOCATION = MaxMemoryAllocation;
constexpr size_t STEP = 1024 * 1024 * 1024ULL;
MemoryStart = nullptr;
for (TotalSize = MAX_MEMORY_ALLOCATION; TotalSize > 0; TotalSize -= STEP) {
auto MemoryStartOrErr =
Device->allocate(TotalSize, /*HstPtr=*/nullptr, TARGET_ALLOC_DEFAULT);
if (!MemoryStartOrErr)
return MemoryStartOrErr.takeError();
MemoryStart = *MemoryStartOrErr;
if (MemoryStart)
break;
}
if (!MemoryStart)
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"allocating record/replay memory");
if (VAddr && VAddr != MemoryStart)
MemoryOffset = uintptr_t(VAddr) - uintptr_t(MemoryStart);
MemoryPtr = MemoryStart;
MemorySize = 0;
// Check if we need adjustment.
if (MemoryOffset > 0 &&
TotalSize >= RequiredMemoryAllocation + MemoryOffset) {
// If we are off but "before" the required address and with enough space,
// we just "allocate" the offset to match the required address.
MemoryPtr = (char *)MemoryPtr + MemoryOffset;
MemorySize += MemoryOffset;
MemoryOffset = 0;
assert(MemoryPtr == VAddr && "Expected offset adjustment to work");
} else if (MemoryOffset) {
// If we are off and in a situation we cannot just "waste" memory to force
// a match, we hope adjusting the arguments is sufficient.
REPORT() << "WARNING Failed to allocate replay memory at required "
<< "location " << VAddr << ", got " << MemoryStart
<< ", trying to offset argument pointers by " << MemoryOffset;
}
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),
"Allocated %" PRIu64 " bytes at %p for replay.\n", TotalSize,
MemoryStart);
return Plugin::success();
}
Error preallocateDeviceMemory(uint64_t DeviceMemorySize, void *ReqVAddr) {
if (Device->supportVAManagement()) {
auto Err = preAllocateVAMemory(DeviceMemorySize, ReqVAddr);
if (Err) {
REPORT() << "WARNING VA mapping failed, fallback to heuristic: "
<< "(Error: " << toString(std::move(Err)) << ")";
}
}
uint64_t DevMemSize;
if (Device->getDeviceMemorySize(DevMemSize))
return Plugin::error(ErrorCode::UNKNOWN,
"cannot determine Device Memory Size");
return preAllocateHeuristic(DevMemSize, DeviceMemorySize, ReqVAddr);
}
void dumpDeviceMemory(StringRef Filename) {
ErrorOr<std::unique_ptr<WritableMemoryBuffer>> DeviceMemoryMB =
WritableMemoryBuffer::getNewUninitMemBuffer(MemorySize);
if (!DeviceMemoryMB)
report_fatal_error("Error creating MemoryBuffer for device memory");
auto Err = Device->dataRetrieve(DeviceMemoryMB.get()->getBufferStart(),
MemoryStart, MemorySize, nullptr);
if (Err)
report_fatal_error("Error retrieving data for target pointer");
StringRef DeviceMemory(DeviceMemoryMB.get()->getBufferStart(), MemorySize);
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
if (EC)
report_fatal_error("Error dumping memory to file " + Filename + " :" +
EC.message());
OS << DeviceMemory;
OS.close();
}
public:
bool isRecording() const { return Status == RRStatusTy::RRRecording; }
bool isReplaying() const { return Status == RRStatusTy::RRReplaying; }
bool isRecordingOrReplaying() const {
return (Status != RRStatusTy::RRDeactivated);
}
void setStatus(RRStatusTy Status) { this->Status = Status; }
bool isSaveOutputEnabled() const { return ReplaySaveOutput; }
void addEntry(const char *Name, uint64_t Size, void *Addr) {
GlobalEntries.emplace_back(GlobalEntry{Name, Size, Addr});
}
void saveImage(const char *Name, const DeviceImageTy &Image) {
SmallString<128> ImageName = {Name, ".image"};
std::error_code EC;
raw_fd_ostream OS(ImageName, EC);
if (EC)
report_fatal_error("Error saving image : " + StringRef(EC.message()));
OS << Image.getMemoryBuffer().getBuffer();
OS.close();
}
void dumpGlobals(StringRef Filename, DeviceImageTy &Image) {
int32_t Size = 0;
for (auto &OffloadEntry : GlobalEntries) {
if (!OffloadEntry.Size)
continue;
// Get the total size of the string and entry including the null byte.
Size += std::strlen(OffloadEntry.Name) + 1 + sizeof(uint32_t) +
OffloadEntry.Size;
}
ErrorOr<std::unique_ptr<WritableMemoryBuffer>> GlobalsMB =
WritableMemoryBuffer::getNewUninitMemBuffer(Size);
if (!GlobalsMB)
report_fatal_error("Error creating MemoryBuffer for globals memory");
void *BufferPtr = GlobalsMB.get()->getBufferStart();
for (auto &OffloadEntry : GlobalEntries) {
if (!OffloadEntry.Size)
continue;
int32_t NameLength = std::strlen(OffloadEntry.Name) + 1;
memcpy(BufferPtr, OffloadEntry.Name, NameLength);
BufferPtr = utils::advancePtr(BufferPtr, NameLength);
*((uint32_t *)(BufferPtr)) = OffloadEntry.Size;
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));
auto Err = Plugin::success();
{
if (auto Err = Device->dataRetrieve(BufferPtr, OffloadEntry.Addr,
OffloadEntry.Size, nullptr))
report_fatal_error("Error retrieving data for global");
}
if (Err)
report_fatal_error("Error retrieving data for global");
BufferPtr = utils::advancePtr(BufferPtr, OffloadEntry.Size);
}
assert(BufferPtr == GlobalsMB->get()->getBufferEnd() &&
"Buffer over/under-filled.");
assert(Size == utils::getPtrDiff(BufferPtr,
GlobalsMB->get()->getBufferStart()) &&
"Buffer size mismatch");
StringRef GlobalsMemory(GlobalsMB.get()->getBufferStart(), Size);
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
OS << GlobalsMemory;
OS.close();
}
void saveKernelDescr(const char *Name, KernelLaunchParamsTy LaunchParams,
int32_t NumArgs, uint64_t NumTeamsClause,
uint32_t ThreadLimitClause, uint64_t LoopTripCount) {
json::Object JsonKernelInfo;
JsonKernelInfo["Name"] = Name;
JsonKernelInfo["NumArgs"] = NumArgs;
JsonKernelInfo["NumTeamsClause"] = NumTeamsClause;
JsonKernelInfo["ThreadLimitClause"] = ThreadLimitClause;
JsonKernelInfo["LoopTripCount"] = LoopTripCount;
JsonKernelInfo["DeviceMemorySize"] = MemorySize;
JsonKernelInfo["DeviceId"] = Device->getDeviceId();
JsonKernelInfo["BumpAllocVAStart"] = (intptr_t)MemoryStart;
json::Array JsonArgPtrs;
for (int I = 0; I < NumArgs; ++I)
JsonArgPtrs.push_back((intptr_t)LaunchParams.Ptrs[I]);
JsonKernelInfo["ArgPtrs"] = json::Value(std::move(JsonArgPtrs));
json::Array JsonArgOffsets;
for (int I = 0; I < NumArgs; ++I)
JsonArgOffsets.push_back(0);
JsonKernelInfo["ArgOffsets"] = json::Value(std::move(JsonArgOffsets));
SmallString<128> JsonFilename = {Name, ".json"};
std::error_code EC;
raw_fd_ostream JsonOS(JsonFilename.str(), EC);
if (EC)
report_fatal_error("Error saving kernel json file : " +
StringRef(EC.message()));
JsonOS << json::Value(std::move(JsonKernelInfo));
JsonOS.close();
}
void saveKernelInput(const char *Name, DeviceImageTy &Image) {
SmallString<128> GlobalsFilename = {Name, ".globals"};
dumpGlobals(GlobalsFilename, Image);
SmallString<128> MemoryFilename = {Name, ".memory"};
dumpDeviceMemory(MemoryFilename);
}
void saveKernelOutputInfo(const char *Name) {
SmallString<128> OutputFilename = {
Name, (isRecording() ? ".original.output" : ".replay.output")};
dumpDeviceMemory(OutputFilename);
}
void *alloc(uint64_t Size) {
assert(MemoryStart && "Expected memory has been pre-allocated");
void *Alloc = nullptr;
constexpr int Alignment = 16;
// Assumes alignment is a power of 2.
int64_t AlignedSize = (Size + (Alignment - 1)) & (~(Alignment - 1));
std::lock_guard<std::mutex> LG(AllocationLock);
Alloc = MemoryPtr;
MemoryPtr = (char *)MemoryPtr + AlignedSize;
MemorySize += AlignedSize;
ODBG(OLDT_Alloc) << "Memory Allocator return " << Alloc;
return Alloc;
}
Error init(GenericDeviceTy *Device, uint64_t MemSize, void *VAddr,
RRStatusTy Status, bool SaveOutput, uint64_t &ReqPtrArgOffset) {
this->Device = Device;
this->Status = Status;
this->ReplaySaveOutput = SaveOutput;
if (auto Err = preallocateDeviceMemory(MemSize, VAddr))
return Err;
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device->getDeviceId(),
"Record Replay Initialized (%p)"
" as starting address, %lu Memory Size"
" and set on status %s\n",
MemoryStart, TotalSize,
Status == RRStatusTy::RRRecording ? "Recording" : "Replaying");
// Tell the user to offset pointer arguments as the memory allocation does
// not match.
ReqPtrArgOffset = MemoryOffset;
return Plugin::success();
}
Error deinit() {
if (UsedVAMap) {
if (auto Err = Device->memoryVAUnMap(MemoryStart, TotalSize))
return Err;
} else {
if (auto Err = Device->free(MemoryStart))
return Err;
}
return Plugin::success();
}
};
} // namespace llvm::omp::target::plugin
AsyncInfoWrapperTy::AsyncInfoWrapperTy(GenericDeviceTy &Device,
__tgt_async_info *AsyncInfoPtr)
: Device(Device),
@ -441,7 +107,8 @@ GenericKernelTy::getKernelLaunchEnvironment(
// Ctor/Dtor have no arguments, replaying uses the original kernel launch
// environment. Older versions of the compiler do not generate a kernel
// launch environment.
if (GenericDevice.Plugin.getRecordReplay().isReplaying() ||
if ((GenericDevice.getRecordReplay() &&
GenericDevice.getRecordReplay()->isReplaying()) ||
KernelArgs.Version < OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR)
return nullptr;
@ -564,7 +231,8 @@ GenericKernelTy::prepareBlockMemory(GenericDeviceTy &GenericDevice,
Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
ptrdiff_t *ArgOffsets, KernelArgsTy &KernelArgs,
AsyncInfoWrapperTy &AsyncInfoWrapper) const {
AsyncInfoWrapperTy &AsyncInfoWrapper,
RecordReplayTy::HandleTy *RRHandle) const {
llvm::SmallVector<void *, 16> Args;
llvm::SmallVector<void *, 16> Ptrs;
@ -608,13 +276,15 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
// Record the kernel description after we modified the argument count and num
// blocks/threads.
RecordReplayTy &RecordReplay = GenericDevice.Plugin.getRecordReplay();
if (RecordReplay.isRecording()) {
RecordReplay.saveImage(getName(), getImage());
RecordReplay.saveKernelInput(getName(), getImage());
RecordReplay.saveKernelDescr(getName(), LaunchParams, KernelArgs.NumArgs,
NumBlocks[0], NumThreads[0],
KernelArgs.Tripcount);
RecordReplayTy *RecordReplay = GenericDevice.getRecordReplay();
if (RecordReplay) {
auto RRHandleOrErr =
RecordReplay->recordPrologue(*this, KernelArgs, LaunchParams, NumBlocks,
NumThreads, DynBlockMemConf.NativeSize);
if (!RRHandleOrErr)
return RRHandleOrErr.takeError();
if (RRHandle)
*RRHandle = *RRHandleOrErr;
}
if (auto Err =
@ -926,10 +596,12 @@ Error GenericDeviceTy::deinit(GenericPluginTy &Plugin) {
delete MemoryManager;
MemoryManager = nullptr;
RecordReplayTy &RecordReplay = Plugin.getRecordReplay();
if (RecordReplay.isRecordingOrReplaying())
if (auto Err = RecordReplay.deinit())
if (RecordReplay) {
if (auto Err = RecordReplay->deinit())
return Err;
delete RecordReplay;
RecordReplay = nullptr;
}
if (RPCServer)
if (auto Err = RPCServer->deinitDevice(*this))
@ -1270,8 +942,8 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
TargetAllocTy Kind) {
void *Alloc = nullptr;
if (Plugin.getRecordReplay().isRecordingOrReplaying())
return Plugin.getRecordReplay().alloc(Size);
if (RecordReplay && RecordReplay->isRecordingOrReplaying())
return RecordReplay->allocate(Size);
switch (Kind) {
case TARGET_ALLOC_DEFAULT:
@ -1335,8 +1007,8 @@ Expected<void *> GenericDeviceTy::dataAlloc(int64_t Size, void *HostPtr,
Error GenericDeviceTy::dataDelete(void *TgtPtr, TargetAllocTy Kind) {
// Free is a noop when recording or replaying.
if (Plugin.getRecordReplay().isRecordingOrReplaying())
return Plugin::success();
if (RecordReplay && RecordReplay->isRecordingOrReplaying())
return RecordReplay->deallocate(TgtPtr);
// Keep track of the deallocation stack if we track allocation traces.
if (OMPX_TrackAllocationTraces) {
@ -1434,9 +1106,8 @@ Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
ptrdiff_t *ArgOffsets,
KernelArgsTy &KernelArgs,
__tgt_async_info *AsyncInfo) {
AsyncInfoWrapperTy AsyncInfoWrapper(
*this,
Plugin.getRecordReplay().isRecordingOrReplaying() ? nullptr : AsyncInfo);
AsyncInfoWrapperTy AsyncInfoWrapper(*this,
RecordReplay ? nullptr : AsyncInfo);
GenericKernelTy &GenericKernel =
*reinterpret_cast<GenericKernelTy *>(EntryPtr);
@ -1453,16 +1124,16 @@ Error GenericDeviceTy::launchKernel(void *EntryPtr, void **ArgPtrs,
.emplace(&GenericKernel, std::move(StackTrace), AsyncInfo);
}
RecordReplayTy::HandleTy RRHandle;
auto Err = GenericKernel.launch(*this, ArgPtrs, ArgOffsets, KernelArgs,
AsyncInfoWrapper);
AsyncInfoWrapper, &RRHandle);
// 'finalize' here to guarantee next record-replay actions are in-sync
AsyncInfoWrapper.finalize(Err);
RecordReplayTy &RecordReplay = Plugin.getRecordReplay();
if (RecordReplay.isRecordingOrReplaying() &&
RecordReplay.isSaveOutputEnabled())
RecordReplay.saveKernelOutputInfo(GenericKernel.getName());
if (RecordReplay)
if (auto Err = RecordReplay->recordEpilogue(GenericKernel, RRHandle))
return Err;
return Err;
}
@ -1605,9 +1276,6 @@ Error GenericPluginTy::init() {
RPCServer = new RPCServerTy(*this);
assert(RPCServer && "Invalid RPC server");
RecordReplay = new RecordReplayTy();
assert(RecordReplay && "Invalid RR interface");
return Plugin::success();
}
@ -1633,9 +1301,6 @@ Error GenericPluginTy::deinit() {
delete RPCServer;
}
if (RecordReplay)
delete RecordReplay;
// Perform last deinitializations on the plugin.
if (Error Err = deinitImpl())
return Err;
@ -1798,25 +1463,19 @@ int32_t GenericPluginTy::is_data_exchangable(int32_t SrcDeviceId,
return isDataExchangable(SrcDeviceId, DstDeviceId);
}
int32_t GenericPluginTy::initialize_record_replay(int32_t DeviceId,
int64_t MemorySize,
void *VAddr, bool isRecord,
bool SaveOutput,
uint64_t &ReqPtrArgOffset) {
int32_t GenericPluginTy::initialize_record_replay(
int32_t DeviceId, int64_t MemorySize, void *VAddr, bool IsRecord,
bool IsNative, bool SaveOutput, bool EmitReport,
const char *OutputDirPath) {
GenericDeviceTy &Device = getDevice(DeviceId);
RecordReplayTy::RRStatusTy Status =
isRecord ? RecordReplayTy::RRStatusTy::RRRecording
: RecordReplayTy::RRStatusTy::RRReplaying;
if (auto Err = RecordReplay->init(&Device, MemorySize, VAddr, Status,
SaveOutput, ReqPtrArgOffset)) {
REPORT() << "WARNING RR did not initialize RR-properly with " << MemorySize
<< " bytes (Error: " << toString(std::move(Err)) << ")";
RecordReplay->setStatus(RecordReplayTy::RRStatusTy::RRDeactivated);
if (!isRecord) {
return OFFLOAD_FAIL;
}
if (auto Err =
Device.initRecordReplay(MemorySize, VAddr, IsRecord, IsNative,
SaveOutput, EmitReport, OutputDirPath)) {
REPORT() << "Failure to initialize RR with " << MemorySize
<< " bytes on device " << DeviceId << ": "
<< toString(std::move(Err));
return OFFLOAD_FAIL;
}
return OFFLOAD_SUCCESS;
}
@ -2185,9 +1844,9 @@ int32_t GenericPluginTy::get_global(__tgt_device_binary Binary, uint64_t Size,
assert(DevicePtr && "Invalid device global's address");
// Save the loaded globals if we are recording.
RecordReplayTy &RecordReplay = Device.Plugin.getRecordReplay();
if (RecordReplay.isRecording())
RecordReplay.addEntry(Name, Size, *DevicePtr);
RecordReplayTy *RecordReplay = Device.getRecordReplay();
if (RecordReplay && RecordReplay->isRecording())
RecordReplay->addGlobal(Name, Size, *DevicePtr);
return OFFLOAD_SUCCESS;
}

328
offload/plugins-nextgen/common/src/RecordReplay.cpp Normal file
View file

@ -0,0 +1,328 @@
//===- RecordReplay.cpp - Target independent kernel record replay ---------===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
//===----------------------------------------------------------------------===//
#include "PluginInterface.h"
#include "Shared/APITypes.h"
#include "ErrorReporting.h"
#include "Shared/Utils.h"
#include "llvm/Support/Error.h"
#include "llvm/Support/JSON.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/raw_ostream.h"
#include <cstdint>
#include <functional>
using namespace llvm;
using namespace omp;
using namespace target;
using namespace plugin;
using namespace error;
RecordReplayTy::InstanceTy::InstanceTy(const GenericKernelTy &Kernel,
uint32_t NumTeams, uint32_t NumThreads,
uint32_t SharedMemorySize)
: Kernel(Kernel), NumTeams(NumTeams), NumThreads(NumThreads),
SharedMemorySize(SharedMemorySize) {
KernelHash = stable_hash_name(Kernel.getName());
LaunchConfigHash =
stable_hash_combine((stable_hash)NumTeams, (stable_hash)NumThreads,
(stable_hash)SharedMemorySize);
}
Error RecordReplayTy::init(uint64_t MemSize, void *VAddr) {
if (!VAddr && isReplaying())
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"VAddr cannot be null when replaying");
if (!VAddr)
VAddr = Device.getSuggestedVirtualAddress();
auto StartAddrOrErr = Device.allocateWithVirtualAddress(MemSize, VAddr);
if (!StartAddrOrErr)
return StartAddrOrErr.takeError();
if (!*StartAddrOrErr)
return Plugin::error(ErrorCode::OUT_OF_RESOURCES, "allocating memory");
if (isReplaying() && *StartAddrOrErr != VAddr)
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"could not reserve recorded virtual address");
StartAddr = *StartAddrOrErr;
TotalSize = MemSize;
// Create the output directory if necessary.
std::error_code EC;
std::filesystem::create_directories(OutputDirectory, EC);
if (EC)
return Plugin::error(ErrorCode::HOST_IO, "creating output directory");
INFO(OMP_INFOTYPE_PLUGIN_KERNEL, Device.getDeviceId(),
"%s initialized with starting address %p, "
"memory size %lu bytes, and output directory in %s\n",
Status == StatusTy::Recording ? "Record" : "Replay", StartAddr,
TotalSize, OutputDirectory.c_str());
return Plugin::success();
}
Error RecordReplayTy::deinit() {
if (isRecording() && EmitReport)
if (auto Err = emitInstanceReport())
return Err;
if (StartAddr)
return Device.deallocateWithVirtualAddress(StartAddr, TotalSize);
return Plugin::success();
}
Error RecordReplayTy::emitInstanceReport() {
std::lock_guard<std::mutex> LG(InstancesLock);
llvm::outs() << "=== record report begin ===\n";
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";
}
llvm::outs() << "=== record report end ===\n";
return Plugin::success();
}
std::pair<const RecordReplayTy::InstanceTy &, bool>
RecordReplayTy::registerInstance(const GenericKernelTy &Kernel,
uint32_t NumTeams, uint32_t NumThreads,
uint32_t SharedMemorySize) {
std::lock_guard<std::mutex> LG(InstancesLock);
auto [It, Inserted] =
Instances.emplace(Kernel, NumTeams, NumThreads, SharedMemorySize);
// Increase the number of occurrences.
It->Occurrences += 1;
return {*It, Inserted};
}
Expected<void *> RecordReplayTy::allocate(uint64_t Size) {
assert(StartAddr && "Expected memory has been pre-allocated");
constexpr int Alignment = 16;
// Assume alignment is a power of 2.
int64_t AlignedSize = (Size + (Alignment - 1)) & (~(Alignment - 1));
std::lock_guard<std::mutex> LG(AllocationLock);
if (CurrentSize + AlignedSize > TotalSize)
return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
"run out of record replay memory");
void *Alloc = (char *)StartAddr + CurrentSize;
CurrentSize += AlignedSize;
return Alloc;
}
Error RecordReplayTy::deallocate(void *Ptr) { return Plugin::success(); }
Expected<RecordReplayTy::HandleTy> RecordReplayTy::recordPrologue(
const GenericKernelTy &Kernel, const KernelArgsTy &KernelArgs,
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);
HandleTy Handle{&Instance, First};
if (isReplaying() || !First)
return Handle;
if (auto Err = recordDescImpl(Kernel, Instance, KernelArgs, LaunchParams))
return Err;
if (auto Err = recordPrologueImpl(Kernel, Instance, KernelArgs, LaunchParams))
return Err;
return Handle;
}
Error RecordReplayTy::recordEpilogue(const GenericKernelTy &Kernel,
HandleTy Handle) {
if (!shouldRecordEpilogue() || !Handle.Active)
return Plugin::success();
return recordEpilogueImpl(Kernel, *Handle.Instance);
}
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))
return Err;
std::string GlobalsFilename = getFilename(Instance, "globals");
if (auto Err = recordGlobals(GlobalsFilename))
return Err;
std::string ImageFilename = getFilename(Instance, "image");
return recordImage(Kernel, ImageFilename);
}
Error NativeRecordReplayTy::recordEpilogueImpl(const GenericKernelTy &Kernel,
const InstanceTy &Instance) {
std::string SnapshotFilename =
getFilename(Instance, isRecording() ? "record_output" : "replay_output");
return recordSnapshot(SnapshotFilename);
}
Error NativeRecordReplayTy::recordDescImpl(
const GenericKernelTy &Kernel, const InstanceTy &Instance,
const KernelArgsTy &KernelArgs, const KernelLaunchParamsTy &LaunchParams) {
json::Object JsonKernelInfo;
JsonKernelInfo["Name"] = Kernel.getName();
JsonKernelInfo["NumArgs"] = KernelArgs.NumArgs;
JsonKernelInfo["NumTeamsClause"] = Instance.NumTeams;
JsonKernelInfo["ThreadLimitClause"] = Instance.NumThreads;
JsonKernelInfo["SharedMemorySize"] = Instance.SharedMemorySize;
JsonKernelInfo["LoopTripCount"] = KernelArgs.Tripcount;
JsonKernelInfo["DeviceId"] = Device.getDeviceId();
JsonKernelInfo["VAllocAddr"] = (intptr_t)StartAddr;
JsonKernelInfo["VAllocSize"] = TotalSize;
json::Array JsonArgPtrs;
for (uint32_t I = 0; I < KernelArgs.NumArgs; ++I)
JsonArgPtrs.push_back((intptr_t)(*(void **)LaunchParams.Ptrs[I]));
JsonKernelInfo["ArgPtrs"] = json::Value(std::move(JsonArgPtrs));
json::Array JsonArgOffsets;
for (uint32_t I = 0; I < KernelArgs.NumArgs; ++I)
JsonArgOffsets.push_back(0);
JsonKernelInfo["ArgOffsets"] = json::Value(std::move(JsonArgOffsets));
std::string JsonFilename = getFilename(Instance, "json");
std::error_code EC;
raw_fd_ostream JsonOS(JsonFilename, EC);
if (EC)
return Plugin::error(ErrorCode::HOST_IO, "saving kernel json file");
JsonOS << json::Value(std::move(JsonKernelInfo));
JsonOS.close();
return Plugin::success();
}
std::string NativeRecordReplayTy::getFilenameImpl(const InstanceTy &Instance,
StringRef Suffix,
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();
}
Error NativeRecordReplayTy::recordSnapshot(const std::string &Filename) {
// Another thread may be allocating memory. The size can only increase.
AllocationLock.lock();
uint64_t RecordSize = CurrentSize;
AllocationLock.unlock();
ErrorOr<std::unique_ptr<WritableMemoryBuffer>> DeviceMemoryMB =
WritableMemoryBuffer::getNewUninitMemBuffer(RecordSize);
if (!DeviceMemoryMB)
return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
"creating MemoryBuffer for device memory");
if (auto Err = Device.dataRetrieve(DeviceMemoryMB.get()->getBufferStart(),
StartAddr, RecordSize, nullptr))
return Err;
StringRef DeviceMemory(DeviceMemoryMB.get()->getBufferStart(), RecordSize);
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
if (EC)
return Plugin::error(ErrorCode::HOST_IO, "saving memory snapshot file");
OS << DeviceMemory;
OS.close();
return Plugin::success();
}
Error NativeRecordReplayTy::recordImage(const GenericKernelTy &Kernel,
const std::string &Filename) {
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
if (EC)
return Plugin::error(ErrorCode::HOST_IO, "saving image file");
OS << Kernel.getImage().getMemoryBuffer().getBuffer();
OS.close();
return Plugin::success();
}
Error NativeRecordReplayTy::recordGlobals(const std::string &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
// the entire function.
SmallVector<GlobalEntryTy> Globals = GlobalEntries;
AllocationLock.unlock();
uint64_t TotalSize = sizeof(uint32_t);
uint32_t NumGlobals = 0;
for (auto &Global : Globals) {
if (!Global.Size)
continue;
// Get the total size of the string and entry including the null byte.
TotalSize += Global.Size + sizeof(uint32_t) + sizeof(uint64_t) +
Global.Name.length() + 1;
NumGlobals++;
}
ErrorOr<std::unique_ptr<WritableMemoryBuffer>> GlobalsMB =
WritableMemoryBuffer::getNewUninitMemBuffer(TotalSize);
if (!GlobalsMB)
return Plugin::error(ErrorCode::OUT_OF_RESOURCES,
"creating MemoryBuffer for globals memory");
void *BufferPtr = GlobalsMB.get()->getBufferStart();
*((uint32_t *)(BufferPtr)) = NumGlobals;
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));
for (auto &Global : Globals) {
if (!Global.Size)
continue;
uint32_t NameLength = Global.Name.length() + 1;
*((uint32_t *)(BufferPtr)) = NameLength;
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));
*((uint64_t *)(BufferPtr)) = Global.Size;
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint64_t));
memcpy(BufferPtr, Global.Name.data(), NameLength);
BufferPtr = utils::advancePtr(BufferPtr, NameLength);
if (auto Err =
Device.dataRetrieve(BufferPtr, Global.Addr, Global.Size, nullptr))
return Err;
BufferPtr = utils::advancePtr(BufferPtr, Global.Size);
}
assert(BufferPtr == GlobalsMB->get()->getBufferEnd() &&
"Buffer over or under-filled.");
assert(TotalSize ==
utils::getPtrDiff(BufferPtr, GlobalsMB->get()->getBufferStart()) &&
"Buffer size mismatch.");
StringRef GlobalsMemory(GlobalsMB.get()->getBufferStart(), TotalSize);
std::error_code EC;
raw_fd_ostream OS(Filename, EC);
OS << GlobalsMemory;
OS.close();
return Plugin::success();
}

179
offload/plugins-nextgen/cuda/src/rtl.cpp
View file

@ -52,38 +52,6 @@ struct CUDAKernelTy;
struct CUDADeviceTy;
struct CUDAPluginTy;
#if (defined(CUDA_VERSION) && (CUDA_VERSION < 11000))
/// Forward declarations for all Virtual Memory Management
/// related data structures and functions. This is necessary
/// for older cuda versions.
typedef void *CUmemGenericAllocationHandle;
typedef void *CUmemAllocationProp;
typedef void *CUmemAccessDesc;
typedef void *CUmemAllocationGranularity_flags;
CUresult cuMemAddressReserve(CUdeviceptr *ptr, size_t size, size_t alignment,
CUdeviceptr addr, unsigned long long flags) {}
CUresult cuMemMap(CUdeviceptr ptr, size_t size, size_t offset,
CUmemGenericAllocationHandle handle,
unsigned long long flags) {}
CUresult cuMemCreate(CUmemGenericAllocationHandle *handle, size_t size,
const CUmemAllocationProp *prop,
unsigned long long flags) {}
CUresult cuMemSetAccess(CUdeviceptr ptr, size_t size,
const CUmemAccessDesc *desc, size_t count) {}
CUresult
cuMemGetAllocationGranularity(size_t *granularity,
const CUmemAllocationProp *prop,
CUmemAllocationGranularity_flags option) {}
#endif
#if (defined(CUDA_VERSION) && (CUDA_VERSION < 11020))
// Forward declarations of asynchronous memory management functions. This is
// necessary for older versions of CUDA.
CUresult cuMemAllocAsync(CUdeviceptr *ptr, size_t, CUstream) { *ptr = 0; }
CUresult cuMemFreeAsync(CUdeviceptr dptr, CUstream hStream) {}
#endif
/// Class implementing the CUDA device images properties.
struct CUDADeviceImageTy : public DeviceImageTy {
/// Create the CUDA image with the id and the target image pointer.
@ -684,125 +652,104 @@ struct CUDADeviceTy : public GenericDeviceTy {
return Plugin::check(Res, "error in cuStreamSynchronize: %s");
}
/// CUDA support VA management
bool supportVAManagement() const override {
#if (defined(CUDA_VERSION) && (CUDA_VERSION >= 11000))
return true;
#else
return false;
#endif
/// Suggest a virtual address for device memory mapping.
void *getSuggestedVirtualAddress() override {
return reinterpret_cast<void *>(0x153940000000ULL);
}
/// Allocates \p RSize bytes (rounded up to page size) and hints the cuda
/// driver to map it to \p VAddr. The obtained address is stored in \p Addr.
/// At return \p RSize contains the actual size
Error memoryVAMap(void **Addr, void *VAddr, size_t *RSize) override {
CUdeviceptr DVAddr = reinterpret_cast<CUdeviceptr>(VAddr);
auto IHandle = DeviceMMaps.find(DVAddr);
size_t Size = *RSize;
/// Allocate \p Size bytes on the device and hints the backend to map it to
/// virtual address \p VAddr. The function returns the allocated virtual
/// address. The memory must be deallocated through
/// GenericDeviceTy::deallocateWithVirtualAddress().
Expected<void *> allocateWithVirtualAddress(uint64_t Size,
void *VAddr) override {
CUdeviceptr ExpectedVAddr = 0;
if (VAddr != nullptr)
ExpectedVAddr = reinterpret_cast<CUdeviceptr>(VAddr);
if (Size == 0)
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"memory Map Size must be larger than 0");
// Check if we have already mapped this address
if (IHandle != DeviceMMaps.end())
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"address already memory mapped");
CUmemAllocationProp Prop = {};
// Get the page size in the device.
size_t Granularity = 0;
size_t Free, Total;
CUresult Res = cuMemGetInfo(&Free, &Total);
if (auto Err = Plugin::check(Res, "Error in cuMemGetInfo: %s"))
return Err;
if (Size >= Free) {
*Addr = nullptr;
return Plugin::error(
ErrorCode::OUT_OF_RESOURCES,
"cannot map memory size larger than the available device memory");
}
// currently NVidia only supports pinned device types
CUmemAllocationProp Prop = {};
Prop.type = CU_MEM_ALLOCATION_TYPE_PINNED;
Prop.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
Prop.location.id = DeviceId;
cuMemGetAllocationGranularity(&Granularity, &Prop,
CU_MEM_ALLOC_GRANULARITY_MINIMUM);
CUresult Res = cuMemGetAllocationGranularity(
&Granularity, &Prop, CU_MEM_ALLOC_GRANULARITY_MINIMUM);
if (auto Err =
Plugin::check(Res, "error in cuMemGetAllocationGranularity: %s"))
return Err;
if (Granularity == 0)
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"wrong device Page size");
"wrong device page size");
// Ceil to page size.
// Transparently round up to a multiple of the page size.
Size = utils::roundUp(Size, Granularity);
// Create a handler of our allocation
CUmemGenericAllocationHandle AHandle;
Res = cuMemCreate(&AHandle, Size, &Prop, 0);
if (auto Err = Plugin::check(Res, "error in cuMemCreate: %s"))
return Err;
// Reserve the virtual address range.
CUdeviceptr DevPtr = 0;
Res = cuMemAddressReserve(&DevPtr, Size, 0, DVAddr, 0);
Res = cuMemAddressReserve(&DevPtr, Size, 0, ExpectedVAddr, 0);
if (auto Err = Plugin::check(Res, "error in cuMemAddressReserve: %s"))
return Err;
Res = cuMemMap(DevPtr, Size, 0, AHandle, 0);
if (ExpectedVAddr != 0 && ExpectedVAddr != DevPtr)
ODBG(OLDT_Alloc) << "cuMemAddressReserve reserved device virtual address "
<< reinterpret_cast<void *>(DevPtr) << " instead of "
<< reinterpret_cast<void *>(ExpectedVAddr);
// Create a handle of the allocation.
CUmemGenericAllocationHandle Handle;
Res = cuMemCreate(&Handle, Size, &Prop, 0);
if (auto Err = Plugin::check(Res, "error in cuMemCreate: %s"))
return Err;
// Map the virtual address range to the memory allocation.
Res = cuMemMap(DevPtr, Size, 0, Handle, 0);
if (auto Err = Plugin::check(Res, "error in cuMemMap: %s"))
return Err;
CUmemAccessDesc ADesc = {};
ADesc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
ADesc.location.id = DeviceId;
ADesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
// Sets address
Res = cuMemSetAccess(DevPtr, Size, &ADesc, 1);
// Set the memory access properties for the allocation.
CUmemAccessDesc Desc = {};
Desc.location.type = CU_MEM_LOCATION_TYPE_DEVICE;
Desc.location.id = DeviceId;
Desc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;
Res = cuMemSetAccess(DevPtr, Size, &Desc, 1);
if (auto Err = Plugin::check(Res, "error in cuMemSetAccess: %s"))
return Err;
*Addr = reinterpret_cast<void *>(DevPtr);
*RSize = Size;
DeviceMMaps.insert({DevPtr, AHandle});
return Plugin::success();
VAddr = reinterpret_cast<void *>(DevPtr);
// Register the virtual address range in the tracker.
if (auto Err = VMemTracker.registerReservation(VAddr, Size, Handle))
return Err;
return VAddr;
}
/// De-allocates device memory and Unmaps the Virtual Addr
Error memoryVAUnMap(void *VAddr, size_t Size) override {
CUdeviceptr DVAddr = reinterpret_cast<CUdeviceptr>(VAddr);
auto IHandle = DeviceMMaps.find(DVAddr);
// Mapping does not exist
if (IHandle == DeviceMMaps.end()) {
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"addr is not MemoryMapped");
}
/// Deallocate device memory \p VAddr, which was allocated through
/// GenericDeviceTy::allocateWithVirtualAddress(), and unmap the virtual
/// address range.
Error deallocateWithVirtualAddress(void *VAddr, uint64_t) override {
// Unregister the virtual address range and obtain the information about
// the reservation.
auto InfoOrErr = VMemTracker.unregisterReservation(VAddr);
if (!InfoOrErr)
return InfoOrErr.takeError();
if (IHandle == DeviceMMaps.end())
return Plugin::error(ErrorCode::INVALID_ARGUMENT,
"addr is not MemoryMapped");
auto [Size, Handle] = *InfoOrErr;
CUdeviceptr DevAddr = reinterpret_cast<CUdeviceptr>(VAddr);
CUmemGenericAllocationHandle &AllocHandle = IHandle->second;
CUresult Res = cuMemUnmap(DVAddr, Size);
CUresult Res = cuMemUnmap(DevAddr, Size);
if (auto Err = Plugin::check(Res, "error in cuMemUnmap: %s"))
return Err;
Res = cuMemRelease(AllocHandle);
Res = cuMemRelease(Handle);
if (auto Err = Plugin::check(Res, "error in cuMemRelease: %s"))
return Err;
Res = cuMemAddressFree(DVAddr, Size);
Res = cuMemAddressFree(DevAddr, Size);
if (auto Err = Plugin::check(Res, "error in cuMemAddressFree: %s"))
return Err;
DeviceMMaps.erase(IHandle);
return Plugin::success();
}
@ -1464,9 +1411,6 @@ private:
/// The CUDA device handler.
CUdevice Device = CU_DEVICE_INVALID;
/// The memory mapped addresses and their handles
std::unordered_map<CUdeviceptr, CUmemGenericAllocationHandle> DeviceMMaps;
/// The compute capability of the corresponding CUDA device.
struct ComputeCapabilityTy {
uint32_t Major;
@ -1479,6 +1423,9 @@ private:
/// The maximum number of warps that can be resident on all the SMs
/// simultaneously.
uint32_t HardwareParallelism = 0;
/// Tracker for virtual address reservations.
VMemTrackerTy<CUmemGenericAllocationHandle> VMemTracker;
};
Error CUDAKernelTy::launchImpl(GenericDeviceTy &GenericDevice,

1
offload/test/lit.cfg
View file

@ -463,3 +463,4 @@ config.substitutions.append(("%not", config.libomptarget_not))
config.substitutions.append(("%offload-device-info",
config.offload_device_info))
config.substitutions.append(("%offload-tblgen", config.offload_tblgen))
config.substitutions.append(("%omp-kernel-replay", config.omp_kernel_replay))

1
offload/test/lit.site.cfg.in
View file

@ -28,5 +28,6 @@ config.libomptarget_debug = @LIBOMPTARGET_DEBUG@
config.has_libomptarget_ompt = @LIBOMPTARGET_OMPT_SUPPORT@
config.libomptarget_has_libc = @LIBOMPTARGET_GPU_LIBC_SUPPORT@
config.offload_tblgen = "@OFFLOAD_TBLGEN_EXECUTABLE@"
config.omp_kernel_replay = "@OMP_KERNEL_REPLAY@"
# Let the main config do the real work.
lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg")

71
offload/test/tools/omp-kernel-replay/record-replay-advanced.cpp Normal file
View file

@ -0,0 +1,71 @@
// 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 '2p' | grep . | xargs -I {} %omp-kernel-replay --verify {}
// RUN: ls -t %t.testdir/*.json | sed -n '3p' | grep . | xargs -I {} %omp-kernel-replay --verify {}
// 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>
#pragma omp declare target
uint64_t GlobalSum = 0;
#pragma omp end declare target
int main() {
size_t Size = 1000;
uint64_t *Data = new uint64_t[Size];
for (size_t I = 0; I < Size; ++I) {
Data[I] = 20;
}
uint64_t Sum = 0;
#pragma omp target data map(tofrom : Data[0 : Size])
{
#pragma omp target teams distribute parallel for num_teams(256) \
thread_limit(128)
for (size_t I = 0; I < Size; ++I) {
Data[I] = 10 + (uint64_t)I;
}
#pragma omp target thread_limit(1) map(tofrom : Sum)
{
for (size_t I = 0; I < Size; ++I) {
GlobalSum += Data[I];
}
Sum = GlobalSum;
}
// CHECK: PASS
if (Sum == 509500)
printf("PASS\n");
#pragma omp target teams distribute parallel for num_teams(512) thread_limit(64)
for (size_t I = 0; I < Size; ++I) {
Data[I] = GlobalSum;
}
}
Sum = 0;
for (size_t I = 0; I < Size; ++I) {
Sum += Data[I];
}
// CHECK: PASS
if (Sum == 509500000ULL)
printf("PASS\n");
delete[] Data;
}

43
offload/test/tools/omp-kernel-replay/record-replay-basic.cpp Normal file
View file

@ -0,0 +1,43 @@
// 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 {}
// 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;
}

153
offload/tools/kernelreplay/llvm-omp-kernel-replay.cpp
View file

@ -11,6 +11,7 @@
//
//===----------------------------------------------------------------------===//
#include "Shared/Utils.h"
#include "omptarget.h"
#include "llvm/Frontend/Offloading/Utility.h"
@ -21,6 +22,7 @@
#include <cstdint>
#include <cstdlib>
#include <filesystem>
using namespace llvm;
@ -42,11 +44,11 @@ static cl::opt<bool> SaveOutputOpt(
cl::desc("Save the device memory output of the replayed kernel execution."),
cl::init(false), cl::cat(ReplayOptions));
static cl::opt<unsigned> NumTeamsOpt("num-teams",
static cl::opt<uint32_t> NumTeamsOpt("num-teams",
cl::desc("Set the number of teams."),
cl::init(0), cl::cat(ReplayOptions));
static cl::opt<unsigned> NumThreadsOpt("num-threads",
static cl::opt<uint32_t> NumThreadsOpt("num-threads",
cl::desc("Set the number of threads."),
cl::init(0), cl::cat(ReplayOptions));
@ -69,16 +71,19 @@ int main(int argc, char **argv) {
auto NumTeamsJson =
JsonKernelInfo->getAsObject()->getInteger("NumTeamsClause");
unsigned NumTeams = (NumTeamsOpt > 0 ? NumTeamsOpt : NumTeamsJson.value());
uint32_t NumTeams = (NumTeamsOpt > 0 ? NumTeamsOpt : NumTeamsJson.value());
auto NumThreadsJson =
JsonKernelInfo->getAsObject()->getInteger("ThreadLimitClause");
unsigned NumThreads =
uint32_t NumThreads =
(NumThreadsOpt > 0 ? NumThreadsOpt : NumThreadsJson.value());
uint32_t SharedMemorySize =
JsonKernelInfo->getAsObject()->getInteger("SharedMemorySize").value();
// TODO: Print a warning if number of teams/threads is explicitly set in the
// kernel info but overridden through command line options.
auto LoopTripCount =
JsonKernelInfo->getAsObject()->getInteger("LoopTripCount");
auto KernelFunc = JsonKernelInfo->getAsObject()->getString("Name");
std::string KernelName = KernelFunc.value().str();
SmallVector<void *> TgtArgs;
SmallVector<ptrdiff_t> TgtArgOffsets;
@ -91,39 +96,77 @@ int main(int argc, char **argv) {
for (auto It : *TgtArgOffsetsArray)
TgtArgOffsets.push_back(static_cast<ptrdiff_t>(It.getAsInteger().value()));
void *BAllocStart = reinterpret_cast<void *>(
JsonKernelInfo->getAsObject()->getInteger("BumpAllocVAStart").value());
void *VAllocAddr = reinterpret_cast<void *>(
JsonKernelInfo->getAsObject()->getInteger("VAllocAddr").value());
uint64_t VAllocSize =
JsonKernelInfo->getAsObject()->getInteger("VAllocSize").value();
llvm::offloading::EntryTy KernelEntry = {~0U, 0, 0, 0, nullptr,
nullptr, 0, 0, nullptr};
std::string KernelEntryName = KernelFunc.value().str();
KernelEntry.SymbolName = const_cast<char *>(KernelEntryName.c_str());
auto Filepath = std::filesystem::path(InputFilename.getValue());
auto Directory = Filepath.parent_path();
Filepath.replace_extension("globals");
ErrorOr<std::unique_ptr<MemoryBuffer>> GlobalsMB =
MemoryBuffer::getFile(Filepath.string(), /*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!GlobalsMB)
reportFatalUsageError("Error reading the globals file");
// On AMD for currently unknown reasons we cannot copy memory mapped data to
// device. This is a work-around.
uint8_t *RecordedGlobals = new uint8_t[GlobalsMB.get()->getBufferSize()];
std::memcpy(RecordedGlobals,
const_cast<char *>(GlobalsMB.get()->getBuffer().data()),
GlobalsMB.get()->getBufferSize());
void *BufferPtr = (void *)RecordedGlobals;
uint32_t NumGlobals = *((uint32_t *)(BufferPtr));
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));
llvm::SmallVector<llvm::offloading::EntryTy> OffloadEntries(
NumGlobals + 1, {0x0, 0x1, object::OffloadKind::OFK_OpenMP, 0, nullptr,
nullptr, 0, 0, nullptr});
OffloadEntries[0].SymbolName = const_cast<char *>(KernelName.c_str());
// Anything non-zero works to uniquely identify the kernel.
KernelEntry.Address = (void *)0x1;
OffloadEntries[0].Address = (void *)0x1;
for (uint32_t I = 0; I < NumGlobals; ++I) {
auto &Global = OffloadEntries[I + 1];
// Use a unique identifier.
Global.Address = static_cast<char *>(OffloadEntries[0].Address) + I + 1;
uint32_t NameSize = *((uint32_t *)(BufferPtr));
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint32_t));
uint64_t Size = *((uint64_t *)(BufferPtr));
BufferPtr = utils::advancePtr(BufferPtr, sizeof(uint64_t));
Global.Size = Size;
Global.SymbolName = (char *)BufferPtr;
BufferPtr = utils::advancePtr(BufferPtr, NameSize);
Global.AuxAddr = BufferPtr;
BufferPtr = utils::advancePtr(BufferPtr, Size);
}
Filepath.replace_extension("image");
ErrorOr<std::unique_ptr<MemoryBuffer>> ImageMB =
MemoryBuffer::getFile(KernelEntryName + ".image", /*isText=*/false,
MemoryBuffer::getFile(Filepath.string(), /*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!ImageMB)
reportFatalUsageError("Error reading the kernel image.");
reportFatalUsageError("Error reading the kernel image file");
__tgt_device_image DeviceImage;
DeviceImage.ImageStart = const_cast<char *>(ImageMB.get()->getBufferStart());
DeviceImage.ImageEnd = const_cast<char *>(ImageMB.get()->getBufferEnd());
DeviceImage.EntriesBegin = &KernelEntry;
DeviceImage.EntriesEnd = &KernelEntry + 1;
DeviceImage.EntriesBegin = &OffloadEntries[0];
DeviceImage.EntriesEnd = &OffloadEntries[OffloadEntries.size() - 1] + 1;
__tgt_bin_desc Desc;
Desc.NumDeviceImages = 1;
Desc.HostEntriesBegin = &KernelEntry;
Desc.HostEntriesEnd = &KernelEntry + 1;
Desc.HostEntriesBegin = &OffloadEntries[0];
Desc.HostEntriesEnd = &OffloadEntries[OffloadEntries.size() - 1] + 1;
Desc.DeviceImages = &DeviceImage;
auto DeviceMemorySizeJson =
JsonKernelInfo->getAsObject()->getInteger("DeviceMemorySize");
// Set device memory size to the ceiling of GB granularity.
uint64_t DeviceMemorySize = std::ceil(DeviceMemorySizeJson.value());
auto DeviceIdJson = JsonKernelInfo->getAsObject()->getInteger("DeviceId");
// TODO: Print warning if the user overrides the device id in the json file.
int32_t DeviceId = (DeviceIdOpt > -1 ? DeviceIdOpt : DeviceIdJson.value());
@ -133,57 +176,51 @@ int main(int argc, char **argv) {
__tgt_register_lib(&Desc);
uint64_t ReqPtrArgOffset = 0;
int Rc = __tgt_activate_record_replay(DeviceId, DeviceMemorySize, BAllocStart,
false, VerifyOpt, ReqPtrArgOffset);
if (Rc != OMP_TGT_SUCCESS) {
report_fatal_error("Cannot activate record replay\n");
}
int Rc = __tgt_activate_record_replay(
DeviceId, VAllocSize, VAllocAddr, /*IsRecord=*/false, VerifyOpt,
/*EmitReport=*/false, Directory.c_str());
if (Rc != OMP_TGT_SUCCESS)
reportFatalUsageError("Error activating record replay");
Filepath.replace_extension("record_input");
ErrorOr<std::unique_ptr<MemoryBuffer>> DeviceMemoryMB =
MemoryBuffer::getFile(KernelEntryName + ".memory", /*isText=*/false,
MemoryBuffer::getFile(Filepath.string(), /*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!DeviceMemoryMB)
reportFatalUsageError("Error reading the kernel input device memory.");
reportFatalUsageError("Error reading the kernel record input file");
// On AMD for currently unknown reasons we cannot copy memory mapped data to
// device. This is a work-around.
uint8_t *recored_data = new uint8_t[DeviceMemoryMB.get()->getBufferSize()];
std::memcpy(recored_data,
uint8_t *RecordedData = new uint8_t[DeviceMemoryMB.get()->getBufferSize()];
std::memcpy(RecordedData,
const_cast<char *>(DeviceMemoryMB.get()->getBuffer().data()),
DeviceMemoryMB.get()->getBufferSize());
// If necessary, adjust pointer arguments.
if (ReqPtrArgOffset) {
for (auto *&Arg : TgtArgs) {
auto ArgInt = uintptr_t(Arg);
// Try to find pointer arguments.
if (ArgInt < uintptr_t(BAllocStart) ||
ArgInt >= uintptr_t(BAllocStart) + DeviceMemorySize)
continue;
Arg = reinterpret_cast<void *>(ArgInt - ReqPtrArgOffset);
}
}
__tgt_target_kernel_replay(
/*Loc=*/nullptr, DeviceId, KernelEntry.Address, (char *)recored_data,
DeviceMemoryMB.get()->getBufferSize(), TgtArgs.data(),
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,
LoopTripCount.value());
SharedMemorySize, LoopTripCount.value());
if (Rc != OMP_TGT_SUCCESS)
reportFatalUsageError("Error replaying kernel");
int ErrorDetected = 0;
if (VerifyOpt) {
Filepath.replace_extension("record_output");
ErrorOr<std::unique_ptr<MemoryBuffer>> OriginalOutputMB =
MemoryBuffer::getFile(KernelEntryName + ".original.output",
MemoryBuffer::getFile(Filepath.string(),
/*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!OriginalOutputMB)
reportFatalUsageError(
"Error reading the kernel original output file, make sure "
"LIBOMPTARGET_SAVE_OUTPUT is set when recording");
"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(KernelEntryName + ".replay.output",
MemoryBuffer::getFile(Filepath.string(),
/*isText=*/false,
/*RequiresNullTerminator=*/false);
if (!ReplayOutputMB)
@ -191,14 +228,16 @@ int main(int argc, char **argv) {
StringRef OriginalOutput = OriginalOutputMB.get()->getBuffer();
StringRef ReplayOutput = ReplayOutputMB.get()->getBuffer();
if (OriginalOutput == ReplayOutput)
if (OriginalOutput == ReplayOutput) {
outs() << "[llvm-omp-kernel-replay] Replay device memory verified!\n";
else
} else {
ErrorDetected = 1;
outs() << "[llvm-omp-kernel-replay] Replay device memory failed to "
"verify!\n";
}
}
delete[] recored_data;
delete[] RecordedData;
return 0;
return ErrorDetected;
}