diff --git a/level_zero/core/source/cmdlist/cmdlist.h b/level_zero/core/source/cmdlist/cmdlist.h index e01ec176e5..a5b03e9ad1 100644 --- a/level_zero/core/source/cmdlist/cmdlist.h +++ b/level_zero/core/source/cmdlist/cmdlist.h @@ -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); diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h index 77a700fce7..fd49f01bbc 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h +++ b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h @@ -232,6 +232,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily::estimateAdditionalSizeAppe return additionalSize; } +template +inline void CommandListCoreFamilyImmediate::setPatchingPreamble(bool value) { + this->cmdQImmediate->setPatchingPreamble(value); +} + } // namespace L0 diff --git a/level_zero/core/test/black_box_tests/zello_graph.cpp b/level_zero/core/test/black_box_tests/zello_graph.cpp index 5778f4790d..1b0b8a7ed9 100644 --- a/level_zero/core/test/black_box_tests/zello_graph.cpp +++ b/level_zero/core/test/black_box_tests/zello_graph.cpp @@ -13,8 +13,6 @@ #include #include -#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(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(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::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::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; } diff --git a/level_zero/core/test/unit_tests/experimental/test_graph.cpp b/level_zero/core/test/unit_tests/experimental/test_graph.cpp index e175208937..46ec9355d9 100644 --- a/level_zero/core/test/unit_tests/experimental/test_graph.cpp +++ b/level_zero/core/test/unit_tests/experimental/test_graph.cpp @@ -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; +using GraphTestInstantiationTest = Test; -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 *>(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(cmdList.begin(), cmdList.end()); + ASSERT_LT(1u, sdiCmds.size()); + + size_t bbStartIdx = 0; + for (auto &sdiCmd : sdiCmds) { + auto storeDataImm = reinterpret_cast(*sdiCmd); + + bbStartDwordBuffer[bbStartIdx] = storeDataImm->getDataDword0(); + bbStartIdx++; + if (storeDataImm->getStoreQword()) { + bbStartDwordBuffer[bbStartIdx] = storeDataImm->getDataDword1(); + bbStartIdx++; + } + } + + MI_BATCH_BUFFER_START *chainBackBbStartCmd = genCmdCast(bbStartDwordBuffer); + ASSERT_NE(nullptr, chainBackBbStartCmd); + + EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListDestroy(immCmdList)); +} + TEST(GraphExecution, GivenEmptyExecutableGraphWhenSubmittingItToCommandListThenTakeCareOnlyOfEvents) { GraphsCleanupGuard graphCleanup; Mock signalEvents[2]; diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp index 7c3761a642..444bac52d0 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp @@ -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) { diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp index 1a138fbc57..6297bd6824 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp @@ -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); diff --git a/level_zero/experimental/source/graph/graph.cpp b/level_zero/experimental/source/graph/graph.cpp index fe61ac4c79..887d1aa05e 100644 --- a/level_zero/experimental/source/graph/graph.cpp +++ b/level_zero/experimental/source/graph/graph.cpp @@ -399,7 +399,9 @@ ze_result_t ExecutableGraph::execute(L0::CommandList *executionTarget, void *pNe auto currSignalEvent = (myLastCommandList == *cmdList) ? hSignalEvent : nullptr; ze_command_list_handle_t hCmdList = *cmdList; + executionTarget->setPatchingPreamble(true); auto res = executionTarget->appendCommandLists(1, &hCmdList, currSignalEvent, numWaitEvents, phWaitEvents); + executionTarget->setPatchingPreamble(false); if (ZE_RESULT_SUCCESS != res) { return res; } @@ -409,7 +411,9 @@ ze_result_t ExecutableGraph::execute(L0::CommandList *executionTarget, void *pNe if (L0::CommandList **cmdList = std::get_if(&this->submissionChain[submissioNodeId])) { auto currSignalEvent = (myLastCommandList == *cmdList) ? hSignalEvent : nullptr; ze_command_list_handle_t hCmdList = *cmdList; + executionTarget->setPatchingPreamble(true); auto res = executionTarget->appendCommandLists(1, &hCmdList, currSignalEvent, 0, nullptr); + executionTarget->setPatchingPreamble(false); if (ZE_RESULT_SUCCESS != res) { return res; }