intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-08 22:12:59 +08:00
Code Issues Packages Projects Releases Wiki Activity

feature: add patch preamble to graphs

Browse Source 
Related-To: NEO-15376

Signed-off-by: Zbigniew Zdanowicz <zbigniew.zdanowicz@intel.com>
This commit is contained in:
Zbigniew Zdanowicz
2025-08-14 09:55:35 +00:00
committed by Compute-Runtime-Automation
parent 259271f59d
commit 1316330ff9
8 changed files with 180 additions and 27 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
level_zero/core/source/cmdlist/cmdlist.h
View File

@@ -486,6 +486,8 @@ struct CommandList : _ze_command_list_handle_t {
return flags;
}
virtual void setPatchingPreamble(bool value) {}
protected:
NEO::GraphicsAllocation *getAllocationFromHostPtrMap(const void *buffer, uint64_t bufferSize, bool copyOffload);
NEO::GraphicsAllocation *getHostPtrAlloc(const void *buffer, uint64_t bufferSize, bool hostCopyAllowed, bool copyOffload);

1
level_zero/core/source/cmdlist/cmdlist_hw_immediate.h
View File

@@ -232,6 +232,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
bool isBarrierRequired();
bool isRelaxedOrderingDispatchAllowed(uint32_t numWaitEvents, bool copyOffload) override;
bool skipInOrderNonWalkerSignalingAllowed(ze_event_handle_t signalEvent) const override;
void setPatchingPreamble(bool value) override;
protected:
using BaseClass::inOrderExecInfo;

5
level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl
View File

@@ -1957,4 +1957,9 @@ size_t CommandListCoreFamilyImmediate<gfxCoreFamily>::estimateAdditionalSizeAppe
return additionalSize;
}
template <GFXCORE_FAMILY gfxCoreFamily>
inline void CommandListCoreFamilyImmediate<gfxCoreFamily>::setPatchingPreamble(bool value) {
this->cmdQImmediate->setPatchingPreamble(value);
}
} // namespace L0

122
level_zero/core/test/black_box_tests/zello_graph.cpp
View File

@@ -13,8 +13,6 @@
#include <cstring>
#include <iostream>
#define ENABLE_GRAPH_DUMP false
using zeGraphCreateExpFP = ze_result_t(ZE_APICALL *)(ze_context_handle_t context, ze_graph_handle_t *phGraph, void *pNext);
using zeCommandListBeginGraphCaptureExpFP = ze_result_t(ZE_APICALL *)(ze_command_list_handle_t hCommandList, void *pNext);
using zeCommandListBeginCaptureIntoGraphExpFP = ze_result_t(ZE_APICALL *)(ze_command_list_handle_t hCommandList, ze_graph_handle_t hGraph, void *pNext);
@@ -49,8 +47,8 @@ struct GraphApi {
bool loaded = false;
};
void dumpGraphToDotIfEnabled(const GraphApi &graphApi, ze_graph_handle_t virtualGraph, const std::string &testName) {
if (!ENABLE_GRAPH_DUMP) {
void dumpGraphToDotIfEnabled(const GraphApi &graphApi, ze_graph_handle_t virtualGraph, const std::string &testName, bool dumpGraph) {
if (!dumpGraph) {
return;
}
@@ -67,13 +65,14 @@ void dumpGraphToDotIfEnabled(const GraphApi &graphApi, ze_graph_handle_t virtual
inline void createImmediateCmdlistWithMode(ze_context_handle_t context,
ze_device_handle_t device,
ze_command_queue_mode_t mode,
ze_command_queue_flags_t flags,
ze_command_list_handle_t &cmdList) {
ze_command_queue_desc_t cmdQueueDesc{
.stype = ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC,
.pNext = nullptr,
.ordinal = LevelZeroBlackBoxTests::getCommandQueueOrdinal(device, false),
.index = 0,
.flags = 0,
.flags = flags,
.mode = mode,
.priority = ZE_COMMAND_QUEUE_PRIORITY_NORMAL,
};
@@ -103,7 +102,7 @@ GraphApi loadGraphApi(ze_driver_handle_t driver) {
return testGraphFunctions;
}
bool testAppendMemoryCopy(GraphApi &graphApi, ze_context_handle_t &context, ze_device_handle_t &device, bool aubMode) {
bool testAppendMemoryCopy(GraphApi &graphApi, ze_context_handle_t &context, ze_device_handle_t &device, bool aubMode, bool dumpGraph) {
bool validRet = true;
ze_graph_handle_t virtualGraph = nullptr;
ze_executable_graph_handle_t physicalGraph = nullptr;
@@ -132,7 +131,7 @@ bool testAppendMemoryCopy(GraphApi &graphApi, ze_context_handle_t &context, ze_d
SUCCESS_OR_TERMINATE(zeMemAllocDevice(context, &deviceDesc, allocSize, allocSize, device, &zeBuffer));
ze_command_list_handle_t cmdList;
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, cmdList);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdList);
SUCCESS_OR_TERMINATE(graphApi.commandListBeginCaptureIntoGraph(cmdList, virtualGraph, nullptr));
@@ -159,7 +158,7 @@ bool testAppendMemoryCopy(GraphApi &graphApi, ze_context_handle_t &context, ze_d
}
}
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__);
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__, dumpGraph);
SUCCESS_OR_TERMINATE(zeMemFree(context, zeBuffer));
SUCCESS_OR_TERMINATE(zeMemFree(context, heapBuffer));
@@ -171,7 +170,7 @@ bool testAppendMemoryCopy(GraphApi &graphApi, ze_context_handle_t &context, ze_d
return validRet;
}
bool testMultiGraph(GraphApi &graphApi, ze_context_handle_t &context, ze_device_handle_t &device, bool aubMode) {
bool testMultiGraph(GraphApi &graphApi, ze_context_handle_t &context, ze_device_handle_t &device, bool aubMode, bool dumpGraph) {
bool validRet = true;
ze_graph_handle_t virtualGraph = nullptr;
ze_executable_graph_handle_t physicalGraph = nullptr;
@@ -219,8 +218,8 @@ bool testMultiGraph(GraphApi &graphApi, ze_context_handle_t &context, ze_device_
eventDesc.index = 1;
SUCCESS_OR_TERMINATE(zeEventCreate(eventPool, &eventDesc, &eventJoin));
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, cmdListMain);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, cmdListSub);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdListMain);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdListSub);
SUCCESS_OR_TERMINATE(graphApi.commandListBeginCaptureIntoGraph(cmdListMain, virtualGraph, nullptr));
@@ -249,7 +248,7 @@ bool testMultiGraph(GraphApi &graphApi, ze_context_handle_t &context, ze_device_
}
}
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__);
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__, dumpGraph);
SUCCESS_OR_TERMINATE(zeMemFree(context, zeBuffer));
SUCCESS_OR_TERMINATE(zeMemFree(context, heapBuffer));
@@ -295,7 +294,8 @@ bool testAppendLaunchKernel(GraphApi &graphApi,
ze_context_handle_t &context,
ze_device_handle_t &device,
bool areDispatchTraitsIndirect,
bool aubMode) {
bool aubMode,
bool dumpGraph) {
bool validRet = true;
ze_module_handle_t module;
LevelZeroBlackBoxTests::createModuleFromSpirV(context, device, LevelZeroBlackBoxTests::memcpyBytesTestKernelSrc, module);
@@ -312,7 +312,7 @@ bool testAppendLaunchKernel(GraphApi &graphApi,
1, &eventCopied, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
ze_command_list_handle_t cmdList;
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, cmdList);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdList);
// Buffers
constexpr size_t allocSize = 4096;
@@ -395,7 +395,7 @@ bool testAppendLaunchKernel(GraphApi &graphApi,
// Validate
if (aubMode == false) {
validRet = LevelZeroBlackBoxTests::validateToValue(srcInitValue, outputData.get(), allocSize);
validRet = LevelZeroBlackBoxTests::validateToValue(srcInitValue, outputData.get(), numThreads);
if (!validRet) {
std::cerr << "Data mismatches found!\n";
std::cerr << "srcInitValue == " << std::dec << srcInitValue << "\n";
@@ -403,7 +403,7 @@ bool testAppendLaunchKernel(GraphApi &graphApi,
}
}
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__);
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__, dumpGraph);
// Cleanup
SUCCESS_OR_TERMINATE(zeMemFree(context, dstBuffer));
@@ -425,7 +425,8 @@ bool testAppendLaunchKernel(GraphApi &graphApi,
bool testAppendLaunchMultipleKernelsIndirect(GraphApi &graphApi,
ze_context_handle_t &context,
ze_device_handle_t &device,
bool aubMode) {
bool aubMode,
bool dumpGraph) {
bool validRet = true;
ze_module_handle_t module;
LevelZeroBlackBoxTests::createModuleFromSpirV(context, device, LevelZeroBlackBoxTests::memcpyBytesAndAddConstTestKernelSrc, module);
@@ -444,7 +445,7 @@ bool testAppendLaunchMultipleKernelsIndirect(GraphApi &graphApi,
1, &eventCopied, ZE_EVENT_SCOPE_FLAG_HOST, ZE_EVENT_SCOPE_FLAG_HOST);
ze_command_list_handle_t cmdList;
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, cmdList);
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdList);
// Buffers
constexpr size_t allocSize = 4096;
@@ -566,7 +567,7 @@ bool testAppendLaunchMultipleKernelsIndirect(GraphApi &graphApi,
validRet &= validRet2; // Combine results
}
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__);
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__, dumpGraph);
// Cleanup
SUCCESS_OR_TERMINATE(zeMemFree(context, dispatchTraits));
@@ -590,17 +591,83 @@ bool testAppendLaunchMultipleKernelsIndirect(GraphApi &graphApi,
return validRet;
}
bool testMultipleGraphExecution(GraphApi &graphApi,
ze_context_handle_t &context,
ze_device_handle_t &device,
bool aubMode,
bool dumpGraph) {
bool validRet = true;
constexpr size_t allocSize = 4096;
constexpr size_t elemCount = allocSize / sizeof(uint32_t);
uint32_t initialValue = 1;
uint32_t loopCount = 3;
uint32_t addValue = 5;
uint32_t expectedValue = initialValue + loopCount * addValue;
ze_module_handle_t module;
LevelZeroBlackBoxTests::createModuleFromSpirV(context, device, LevelZeroBlackBoxTests::memcpyBytesAndAddConstTestKernelSrc, module);
ze_kernel_handle_t kernelAddConstant;
LevelZeroBlackBoxTests::createKernelWithName(module, "add_constant", kernelAddConstant);
ze_command_list_handle_t cmdList;
createImmediateCmdlistWithMode(context, device, ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS, 0, cmdList);
void *buffer = nullptr;
ze_host_mem_alloc_desc_t hostDesc = {ZE_STRUCTURE_TYPE_HOST_MEM_ALLOC_DESC};
SUCCESS_OR_TERMINATE(zeMemAllocHost(context, &hostDesc, allocSize, allocSize, &buffer));
for (size_t i = 0; i < elemCount; i++) {
reinterpret_cast<uint32_t *>(buffer)[i] = initialValue;
}
ze_graph_handle_t virtualGraph = nullptr;
SUCCESS_OR_TERMINATE(graphApi.graphCreate(context, &virtualGraph, nullptr));
SUCCESS_OR_TERMINATE(graphApi.commandListBeginCaptureIntoGraph(cmdList, virtualGraph, nullptr));
uint32_t groupSizeX = 64;
uint32_t groupSizeY = 1u;
uint32_t groupSizeZ = 1u;
SUCCESS_OR_TERMINATE(zeKernelSetGroupSize(kernelAddConstant, groupSizeX, groupSizeY, groupSizeZ));
SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernelAddConstant, 0, sizeof(buffer), &buffer));
SUCCESS_OR_TERMINATE(zeKernelSetArgumentValue(kernelAddConstant, 1, sizeof(addValue), &addValue));
ze_group_count_t groupCount = {static_cast<uint32_t>(elemCount / groupSizeX), 1, 1};
SUCCESS_OR_TERMINATE(zeCommandListAppendLaunchKernel(cmdList, kernelAddConstant, &groupCount, nullptr, 0, nullptr));
SUCCESS_OR_TERMINATE(graphApi.commandListEndGraphCapture(cmdList, nullptr, nullptr));
ze_executable_graph_handle_t physicalGraph = nullptr;
SUCCESS_OR_TERMINATE(graphApi.commandListInstantiateGraph(virtualGraph, &physicalGraph, nullptr));
// Dispatch and wait
for (uint32_t i = 0; i < loopCount; i++) {
SUCCESS_OR_TERMINATE(graphApi.commandListAppendGraph(cmdList, physicalGraph, nullptr, nullptr, 0, nullptr));
}
SUCCESS_OR_TERMINATE(zeCommandListHostSynchronize(cmdList, std::numeric_limits<uint64_t>::max()));
if (aubMode == false) {
validRet = LevelZeroBlackBoxTests::validateToValue(expectedValue, buffer, elemCount);
}
dumpGraphToDotIfEnabled(graphApi, virtualGraph, __func__, dumpGraph);
SUCCESS_OR_TERMINATE(zeMemFree(context, buffer));
SUCCESS_OR_TERMINATE(zeCommandListDestroy(cmdList));
SUCCESS_OR_TERMINATE(zeKernelDestroy(kernelAddConstant));
SUCCESS_OR_TERMINATE(zeModuleDestroy(module));
return validRet;
}
int main(int argc, char *argv[]) {
constexpr uint32_t bitNumberTestStandardMemoryCopy = 0u;
constexpr uint32_t bitNumberTestStandardMemoryCopyMultigraph = 1u;
constexpr uint32_t bitNumberTestAppendLaunchKernel = 2u;
constexpr uint32_t bitNumberTestAppendLaunchKernelIndirect = 3u;
constexpr uint32_t bitNumberTestAppendLaunchMultipleKernelsIndirect = 4u;
constexpr uint32_t bitNumberTestMultipleExecution = 5u;
constexpr uint32_t defaultTestMask = std::numeric_limits<uint32_t>::max();
LevelZeroBlackBoxTests::TestBitMask testMask = LevelZeroBlackBoxTests::getTestMask(argc, argv, defaultTestMask);
LevelZeroBlackBoxTests::verbose = LevelZeroBlackBoxTests::isVerbose(argc, argv);
bool aubMode = LevelZeroBlackBoxTests::isAubMode(argc, argv);
bool dumpGraph = LevelZeroBlackBoxTests::isParamEnabled(argc, argv, "-d", "--dump_graph");
const std::string blackBoxName("Zello Graph");
@@ -625,35 +692,42 @@ int main(int argc, char *argv[]) {
std::string currentTest;
if (testMask.test(bitNumberTestStandardMemoryCopy)) {
currentTest = "Standard Memory Copy";
casePass = testAppendMemoryCopy(graphApi, context, device0, aubMode);
casePass = testAppendMemoryCopy(graphApi, context, device0, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}
if (testMask.test(bitNumberTestStandardMemoryCopyMultigraph)) {
currentTest = "Standard Memory Copy - multigraph";
casePass = testMultiGraph(graphApi, context, device0, aubMode);
casePass = testMultiGraph(graphApi, context, device0, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}
if (testMask.test(bitNumberTestAppendLaunchKernel)) {
currentTest = "AppendLaunchKernel";
casePass = testAppendLaunchKernel(graphApi, context, device0, false, aubMode);
casePass = testAppendLaunchKernel(graphApi, context, device0, false, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}
if (testMask.test(bitNumberTestAppendLaunchKernelIndirect)) {
currentTest = "AppendLaunchKernelIndirect";
casePass = testAppendLaunchKernel(graphApi, context, device0, true, aubMode);
casePass = testAppendLaunchKernel(graphApi, context, device0, true, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}
if (testMask.test(bitNumberTestAppendLaunchMultipleKernelsIndirect)) {
currentTest = "AppendLaunchMultipleKernelsIndirect";
casePass = testAppendLaunchMultipleKernelsIndirect(graphApi, context, device0, aubMode);
casePass = testAppendLaunchMultipleKernelsIndirect(graphApi, context, device0, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}
if (testMask.test(bitNumberTestMultipleExecution)) {
currentTest = "Multiple Graph Execution";
casePass = testMultipleGraphExecution(graphApi, context, device0, aubMode, dumpGraph);
LevelZeroBlackBoxTests::printResult(aubMode, casePass, blackBoxName, currentTest);
boxPass &= casePass;
}

70
level_zero/core/test/unit_tests/experimental/test_graph.cpp
View File

@@ -7,6 +7,12 @@
#include "level_zero/core/test/unit_tests/experimental/test_graph.h"
#include "shared/source/command_container/cmdcontainer.h"
#include "shared/source/command_stream/linear_stream.h"
#include "shared/test/common/cmd_parse/gen_cmd_parse.h"
#include "shared/test/common/test_macros/hw_test.h"
#include "level_zero/core/source/cmdlist/cmdlist_hw_immediate.h"
#include "level_zero/core/test/unit_tests/fixtures/device_fixture.h"
#include "level_zero/core/test/unit_tests/mocks/mock_module.h"
@@ -708,9 +714,9 @@ TEST(GraphInstantiationValidation, WhenSubGraphsAreNotValidForInstantiationThenW
}
}
using GraphTestInstantiationFixture = Test<DeviceFixture>;
using GraphTestInstantiationTest = Test<DeviceFixture>;
TEST_F(GraphTestInstantiationFixture, WhenInstantiatingGraphThenBakeCommandsIntoCommandlists) {
TEST_F(GraphTestInstantiationTest, WhenInstantiatingGraphThenBakeCommandsIntoCommandlists) {
GraphsCleanupGuard graphCleanup;
MockGraphContextReturningSpecificCmdList ctx;
@@ -835,6 +841,66 @@ TEST_F(GraphTestInstantiationFixture, WhenInstantiatingGraphThenBakeCommandsInto
EXPECT_EQ(1U, graphHwCommands->appendLaunchMultipleKernelsIndirectCalled);
}
HWCMDTEST_F(IGFX_XE_HP_CORE,
GraphTestInstantiationTest,
GivenExecutableGraphWhenSubmittingItToCommandListThenPatchPreambleIsUsed) {
using MI_STORE_DATA_IMM = typename FamilyType::MI_STORE_DATA_IMM;
using MI_BATCH_BUFFER_START = typename FamilyType::MI_BATCH_BUFFER_START;
uint32_t bbStartDwordBuffer[alignUp(sizeof(MI_BATCH_BUFFER_START) / sizeof(uint32_t), 2)] = {0};
GraphsCleanupGuard graphCleanup;
ze_command_queue_desc_t cmdQueueDesc = {ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC};
ze_command_list_handle_t immCmdList;
EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListCreateImmediate(context, device, &cmdQueueDesc, &immCmdList));
auto immCmdListHw = static_cast<CommandListCoreFamilyImmediate<FamilyType::gfxCoreFamily> *>(immCmdList);
Graph srcGraph(context, true);
immCmdListHw->setCaptureTarget(&srcGraph);
srcGraph.startCapturingFrom(*immCmdListHw, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, ::zeCommandListAppendBarrier(immCmdList, nullptr, 0U, nullptr));
srcGraph.stopCapturing();
immCmdListHw->setCaptureTarget(nullptr);
ExecutableGraph execGraph;
execGraph.instantiateFrom(srcGraph);
auto cmdStream = immCmdListHw->getCmdContainer().getCommandStream();
void *immListCpuBase = cmdStream->getCpuBase();
auto usedSpaceBefore = cmdStream->getUsed();
auto res = execGraph.execute(immCmdListHw, nullptr, nullptr, 0, nullptr);
EXPECT_EQ(ZE_RESULT_SUCCESS, res);
auto usedSpaceAfter = cmdStream->getUsed();
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::Parse::parseCommandBuffer(
cmdList,
ptrOffset(immListCpuBase, usedSpaceBefore),
usedSpaceAfter - usedSpaceBefore));
auto sdiCmds = findAll<MI_STORE_DATA_IMM *>(cmdList.begin(), cmdList.end());
ASSERT_LT(1u, sdiCmds.size());
size_t bbStartIdx = 0;
for (auto &sdiCmd : sdiCmds) {
auto storeDataImm = reinterpret_cast<MI_STORE_DATA_IMM *>(*sdiCmd);
bbStartDwordBuffer[bbStartIdx] = storeDataImm->getDataDword0();
bbStartIdx++;
if (storeDataImm->getStoreQword()) {
bbStartDwordBuffer[bbStartIdx] = storeDataImm->getDataDword1();
bbStartIdx++;
}
}
MI_BATCH_BUFFER_START *chainBackBbStartCmd = genCmdCast<MI_BATCH_BUFFER_START *>(bbStartDwordBuffer);
ASSERT_NE(nullptr, chainBackBbStartCmd);
EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListDestroy(immCmdList));
}
TEST(GraphExecution, GivenEmptyExecutableGraphWhenSubmittingItToCommandListThenTakeCareOnlyOfEvents) {
GraphsCleanupGuard graphCleanup;
Mock<Event> signalEvents[2];

1
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp
View File

@@ -2543,6 +2543,7 @@ TEST_F(CommandListCreateTests, givenCreatingRegularCommandlistAndppendCommandLis
EXPECT_FALSE(commandList->isImmediateType());
auto result = commandList->appendCommandLists(0u, nullptr, nullptr, 0u, nullptr);
EXPECT_EQ(result, ZE_RESULT_ERROR_INVALID_ARGUMENT);
commandList->setPatchingPreamble(true);
}
HWTEST_F(CommandListCreateTests, GivenGpuHangWhenCreatingImmediateCommandListAndAppendingEventResetThenDeviceLostIsReturned) {

2
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp
View File

@@ -1723,7 +1723,7 @@ HWCMDTEST_F(IGFX_XE_HP_CORE,
uint64_t queueGpuBase = cmdStream->getGpuBase();
auto commandListHandle = commandList->toHandle();
commandListImmediate->cmdQImmediate->setPatchingPreamble(true);
commandListImmediate->setPatchingPreamble(true);
auto usedSpaceBefore = cmdStream->getUsed();
returnValue = commandListImmediate->appendCommandLists(1, &commandListHandle, nullptr, 0, nullptr);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);