intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-30 18:37:46 +08:00
Code Issues Packages Projects Releases Wiki Activity

feature: Record&Replay support for appendLaunchKernel

Browse Source 
Related-To: NEO-15374
Signed-off-by: Maciej Bielski <maciej.bielski@intel.com>
This commit is contained in:
Maciej Bielski
2025-07-25 15:36:04 +00:00
committed by Compute-Runtime-Automation
parent e396b7e5b8
commit c396367411
5 changed files with 239 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

21
level_zero/experimental/source/graph/graph.cpp
View File

@@ -10,6 +10,7 @@
#include "level_zero/core/source/cmdlist/cmdlist.h"
#include "level_zero/core/source/context/context.h"
#include "level_zero/core/source/event/event.h"
#include "level_zero/core/source/kernel/kernel_imp.h"
namespace L0 {
@@ -219,6 +220,26 @@ ze_result_t Closure<CaptureApi::zeCommandListAppendImageCopyFromMemoryExt>::inst
apiArgs.hSignalEvent, apiArgs.numWaitEvents, const_cast<ze_event_handle_t *>(getOptionalData(indirectArgs.waitEvents)));
}
Closure<CaptureApi::zeCommandListAppendLaunchKernel>::Closure(const ApiArgs &apiArgs) : apiArgs{apiArgs} {
this->apiArgs.launchKernelArgs = nullptr;
this->apiArgs.phWaitEvents = nullptr;
this->indirectArgs.launchKernelArgs = *apiArgs.launchKernelArgs;
auto kernel = static_cast<KernelImp *>(Kernel::fromHandle(apiArgs.kernelHandle));
this->indirectArgs.kernelState = kernel->getMutableState();
this->indirectArgs.waitEvents.reserve(apiArgs.numWaitEvents);
for (uint32_t i = 0; i < apiArgs.numWaitEvents; ++i) {
this->indirectArgs.waitEvents.push_back(apiArgs.phWaitEvents[i]);
}
}
ze_result_t Closure<CaptureApi::zeCommandListAppendLaunchKernel>::instantiateTo(L0::CommandList &executionTarget) const {
auto kernel = static_cast<KernelImp *>(Kernel::fromHandle(apiArgs.kernelHandle));
kernel->getMutableState() = this->indirectArgs.kernelState;
return zeCommandListAppendLaunchKernel(&executionTarget, apiArgs.kernelHandle, &indirectArgs.launchKernelArgs, apiArgs.hSignalEvent, apiArgs.numWaitEvents, const_cast<ze_event_handle_t *>(getOptionalData(indirectArgs.waitEvents)));
}
ExecutableGraph::~ExecutableGraph() = default;
L0::CommandList *ExecutableGraph::allocateAndAddCommandListSubmissionNode() {