From 4e8600d8d00e7869e66e37c36ddd723a98b8d01d Mon Sep 17 00:00:00 2001 From: "Dunajski, Bartosz" Date: Tue, 26 Sep 2023 12:23:14 +0000 Subject: [PATCH] feature: initial support for RelaxedOrdering of in-order Events chaining Disabled by default. Related-To: NEO-7966 Signed-off-by: Dunajski, Bartosz --- .../api/core/ze_module_api_entrypoints.h | 7 +- level_zero/core/source/cmdlist/cmdlist.h | 3 + level_zero/core/source/cmdlist/cmdlist_hw.h | 3 + level_zero/core/source/cmdlist/cmdlist_hw.inl | 2 +- .../source/cmdlist/cmdlist_hw_immediate.h | 4 +- .../source/cmdlist/cmdlist_hw_immediate.inl | 33 +++- .../cmdlist/cmdlist_hw_skl_to_tgllp.inl | 6 + .../cmdlist/cmdlist_hw_xehp_and_later.inl | 19 ++- .../core/test/unit_tests/mocks/mock_cmdlist.h | 1 + .../test_cmdlist_append_launch_kernel_3.cpp | 141 +++++++++++++++++- .../debug_settings/debug_variables_base.inl | 1 + shared/test/common/test_files/igdrcl.config | 1 + 12 files changed, 207 insertions(+), 14 deletions(-) diff --git a/level_zero/api/core/ze_module_api_entrypoints.h b/level_zero/api/core/ze_module_api_entrypoints.h index 0895d5993c..7160a76e2e 100644 --- a/level_zero/api/core/ze_module_api_entrypoints.h +++ b/level_zero/api/core/ze_module_api_entrypoints.h @@ -146,8 +146,13 @@ ze_result_t zeCommandListAppendLaunchKernel( ze_event_handle_t hSignalEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) { + + auto cmdList = L0::CommandList::fromHandle(hCommandList); + L0::CmdListKernelLaunchParams launchParams = {}; - return L0::CommandList::fromHandle(hCommandList)->appendLaunchKernel(kernelHandle, launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, launchParams, false); + launchParams.skipInOrderNonWalkerSignaling = cmdList->skipInOrderNonWalkerSignalingAllowed(hSignalEvent); + + return cmdList->appendLaunchKernel(kernelHandle, launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, launchParams, false); } ze_result_t zeCommandListAppendLaunchCooperativeKernel( diff --git a/level_zero/core/source/cmdlist/cmdlist.h b/level_zero/core/source/cmdlist/cmdlist.h index a2eebfecf5..4ab81ae8be 100644 --- a/level_zero/core/source/cmdlist/cmdlist.h +++ b/level_zero/core/source/cmdlist/cmdlist.h @@ -42,6 +42,7 @@ struct CmdListKernelLaunchParams { bool isBuiltInKernel = false; bool isDestinationAllocationInSystemMemory = false; bool isHostSignalScopeEvent = false; + bool skipInOrderNonWalkerSignaling = false; }; struct CmdListReturnPoint { @@ -345,6 +346,8 @@ struct CommandList : _ze_command_list_handle_t { return kernelWithAssertAppended; } + virtual bool skipInOrderNonWalkerSignalingAllowed(ze_event_handle_t signalEvent) const { return false; } + protected: NEO::GraphicsAllocation *getAllocationFromHostPtrMap(const void *buffer, uint64_t bufferSize); NEO::GraphicsAllocation *getHostPtrAlloc(const void *buffer, uint64_t bufferSize, bool hostCopyAllowed); diff --git a/level_zero/core/source/cmdlist/cmdlist_hw.h b/level_zero/core/source/cmdlist/cmdlist_hw.h index 87f1e4f2ce..f586faab8c 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw.h +++ b/level_zero/core/source/cmdlist/cmdlist_hw.h @@ -68,6 +68,8 @@ template struct CommandListCoreFamily : CommandListImp { using GfxFamily = typename NEO::GfxFamilyMapper::GfxFamily; + using CommandListImp::skipInOrderNonWalkerSignalingAllowed; + using CommandListImp::CommandListImp; ze_result_t initialize(Device *device, NEO::EngineGroupType engineGroupType, ze_command_list_flags_t flags) override; void programL3(bool isSLMused); @@ -331,6 +333,7 @@ struct CommandListCoreFamily : CommandListImp { void handleInOrderImplicitDependencies(bool relaxedOrderingAllowed); virtual void handleInOrderDependencyCounter(); bool isQwordInOrderCounter() const { return GfxFamily::isQwordInOrderCounter; } + bool isInOrderNonWalkerSignalingRequired(const Event *event) const; void addCmdForPatching(void *cmd, uint64_t counterValue, InOrderPatchCommandTypes::CmdType cmdType); diff --git a/level_zero/core/source/cmdlist/cmdlist_hw.inl b/level_zero/core/source/cmdlist/cmdlist_hw.inl index 029c5dfd61..401c50ba37 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw.inl @@ -354,7 +354,7 @@ ze_result_t CommandListCoreFamily::appendLaunchKernel(ze_kernel_h auto res = appendLaunchKernelWithParams(Kernel::fromHandle(kernelHandle), threadGroupDimensions, event, launchParams); - if (isInOrderExecutionEnabled()) { + if (isInOrderExecutionEnabled() && !launchParams.skipInOrderNonWalkerSignaling) { handleInOrderDependencyCounter(); } diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h index 749ac24bb3..fdf38e8b09 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h +++ b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h @@ -165,7 +165,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily::hasStallingCmdsForRelaxedOrd return (!relaxedOrderingDispatch && (numWaitEvents > 0 || this->inOrderDependencyCounter > 0)); } +template +bool CommandListCoreFamilyImmediate::skipInOrderNonWalkerSignalingAllowed(ze_event_handle_t signalEvent) const { + return this->isInOrderNonWalkerSignalingRequired(Event::fromHandle(signalEvent)); +} + template ze_result_t CommandListCoreFamilyImmediate::appendLaunchKernel( ze_kernel_handle_t kernelHandle, const ze_group_count_t *threadGroupDimensions, @@ -414,6 +419,7 @@ ze_result_t CommandListCoreFamilyImmediate::appendLaunchKernel( const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch) { relaxedOrderingDispatch = isRelaxedOrderingDispatchAllowed(numWaitEvents); + bool stallingCmdsForRelaxedOrdering = hasStallingCmdsForRelaxedOrdering(numWaitEvents, relaxedOrderingDispatch); checkAvailableSpace(numWaitEvents, relaxedOrderingDispatch, commonImmediateCommandSize); bool hostWait = waitForEventsFromHost(); @@ -429,7 +435,32 @@ ze_result_t CommandListCoreFamilyImmediate::appendLaunchKernel( hSignalEvent, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); - return flushImmediate(ret, true, hasStallingCmdsForRelaxedOrdering(numWaitEvents, relaxedOrderingDispatch), relaxedOrderingDispatch, true, hSignalEvent); + if (isInOrderExecutionEnabled() && launchParams.skipInOrderNonWalkerSignaling) { + // skip only in base appendLaunchKernel() + handleInOrderNonWalkerSignaling(Event::fromHandle(hSignalEvent), stallingCmdsForRelaxedOrdering, relaxedOrderingDispatch, ret); + CommandListCoreFamily::handleInOrderDependencyCounter(); + } + + return flushImmediate(ret, true, stallingCmdsForRelaxedOrdering, relaxedOrderingDispatch, true, hSignalEvent); +} + +template +void CommandListCoreFamilyImmediate::handleInOrderNonWalkerSignaling(Event *event, bool &hasStallingCmds, bool &relaxedOrderingDispatch, ze_result_t &result) { + bool nonWalkerSignalingHasRelaxedOrdering = false; + + if (NEO::DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.get() == 1) { + nonWalkerSignalingHasRelaxedOrdering = isRelaxedOrderingDispatchAllowed(1); + } + + if (nonWalkerSignalingHasRelaxedOrdering) { + result = flushImmediate(result, true, hasStallingCmds, relaxedOrderingDispatch, true, nullptr); + NEO::RelaxedOrderingHelper::encodeRegistersBeforeDependencyCheckers(*this->commandContainer.getCommandStream()); + relaxedOrderingDispatch = true; + hasStallingCmds = hasStallingCmdsForRelaxedOrdering(1, relaxedOrderingDispatch); + } + + CommandListCoreFamily::appendWaitOnSingleEvent(event, nonWalkerSignalingHasRelaxedOrdering); + CommandListCoreFamily::appendSignalInOrderDependencyCounter(); } template diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl b/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl index 41b30cd99f..0a7e1b7313 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl @@ -40,12 +40,18 @@ size_t CommandListCoreFamily::getReserveSshSize() { template void CommandListCoreFamily::adjustWriteKernelTimestamp(uint64_t globalAddress, uint64_t contextAddress, bool maskLsb, uint32_t mask, bool workloadPartition) {} +template +bool CommandListCoreFamily::isInOrderNonWalkerSignalingRequired(const Event *event) const { + return false; +} + template ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(Kernel *kernel, const ze_group_count_t *threadGroupDimensions, Event *event, const CmdListKernelLaunchParams &launchParams) { UNRECOVERABLE_IF(kernel == nullptr); + UNRECOVERABLE_IF(launchParams.skipInOrderNonWalkerSignaling); const auto driverHandle = static_cast(device->getDriverHandle()); const auto &kernelDescriptor = kernel->getKernelDescriptor(); if (kernelDescriptor.kernelAttributes.flags.isInvalid) { diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl b/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl index b655651842..3b4bf784cd 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl @@ -77,9 +77,12 @@ void programEventL3Flush(Event *event, } template -ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(Kernel *kernel, - const ze_group_count_t *threadGroupDimensions, - Event *event, +bool CommandListCoreFamily::isInOrderNonWalkerSignalingRequired(const Event *event) const { + return (event && (event->isUsingContextEndOffset() || !event->isInOrderExecEvent())); +} + +template +ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(Kernel *kernel, const ze_group_count_t *threadGroupDimensions, Event *event, const CmdListKernelLaunchParams &launchParams) { if (NEO::DebugManager.flags.ForcePipeControlPriorToWalker.get()) { @@ -173,7 +176,6 @@ ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(K uint64_t eventAddress = 0; bool isTimestampEvent = false; - bool isInOrderExecEvent = false; bool l3FlushEnable = false; bool isHostSignalScopeEvent = launchParams.isHostSignalScopeEvent; Event *compactEvent = nullptr; @@ -182,7 +184,6 @@ ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(K event->setKernelForPrintf(kernel); } isHostSignalScopeEvent = event->isSignalScope(ZE_EVENT_SCOPE_FLAG_HOST); - isInOrderExecEvent = event->isInOrderExecEvent(); if (compactL3FlushEvent(getDcFlushRequired(event->isSignalScope()))) { compactEvent = event; event = nullptr; @@ -298,7 +299,7 @@ ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(K }; bool inOrderExecSignalRequired = (this->inOrderExecutionEnabled && !launchParams.isKernelSplitOperation); - bool inOrderNonWalkerSignalling = event && (isTimestampEvent || !isInOrderExecEvent); + bool inOrderNonWalkerSignalling = isInOrderNonWalkerSignalingRequired(event); if (inOrderExecSignalRequired) { if (inOrderNonWalkerSignalling) { @@ -329,8 +330,10 @@ ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(K if (inOrderExecSignalRequired) { if (inOrderNonWalkerSignalling) { - appendWaitOnSingleEvent(event, false); - appendSignalInOrderDependencyCounter(); + if (!launchParams.skipInOrderNonWalkerSignaling) { + appendWaitOnSingleEvent(event, false); + appendSignalInOrderDependencyCounter(); + } } else { UNRECOVERABLE_IF(!dispatchKernelArgs.outWalkerPtr); addCmdForPatching(dispatchKernelArgs.outWalkerPtr, dispatchKernelArgs.postSyncImmValue, InOrderPatchCommandTypes::CmdType::Walker); diff --git a/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h b/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h index 9537c93881..d7b143630a 100644 --- a/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h +++ b/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h @@ -176,6 +176,7 @@ struct WhiteBox> using BaseClass::inOrderPatchCmds; using BaseClass::isBcsSplitNeeded; using BaseClass::isFlushTaskSubmissionEnabled; + using BaseClass::isInOrderNonWalkerSignalingRequired; using BaseClass::isQwordInOrderCounter; using BaseClass::isSyncModeQueue; using BaseClass::isTbxMode; diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp index a870ff0dea..7f49258007 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp @@ -22,8 +22,10 @@ #include "shared/test/common/cmd_parse/gen_cmd_parse.h" #include "shared/test/common/helpers/debug_manager_state_restore.h" #include "shared/test/common/helpers/engine_descriptor_helper.h" +#include "shared/test/common/helpers/relaxed_ordering_commands_helper.h" #include "shared/test/common/helpers/unit_test_helper.h" #include "shared/test/common/libult/ult_command_stream_receiver.h" +#include "shared/test/common/mocks/mock_direct_submission_hw.h" #include "shared/test/common/mocks/mock_os_context.h" #include "shared/test/common/test_macros/hw_test.h" @@ -712,7 +714,12 @@ struct InOrderCmdListTests : public CommandListAppendLaunchKernel { template DestroyableZeUniquePtr>> createImmCmdList() { - auto cmdList = makeZeUniquePtr>>(); + return createImmCmdListImpl>>(); + } + + template + DestroyableZeUniquePtr createImmCmdListImpl() { + auto cmdList = makeZeUniquePtr(); auto csr = device->getNEODevice()->getDefaultEngine().commandStreamReceiver; @@ -1517,7 +1524,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingTimestampEventThen auto eventPool = createEvents(1, true); events[0]->signalScope = 0; - immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false); + zeCommandListAppendLaunchKernel(immCmdList->toHandle(), kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr); GenCmdList cmdList; ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, cmdStream->getCpuBase(), cmdStream->getUsed())); @@ -1559,6 +1566,123 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingTimestampEventThen EXPECT_EQ(1u, sdiCmd->getDataDword0()); } +HWTEST2_F(InOrderCmdListTests, givenRelaxedOrderingWhenProgrammingTimestampEventThenClearAndChainWithSyncAllocSignalingAsTwoSeparateSubmissions, IsAtLeastXeHpcCore) { + using MI_STORE_DATA_IMM = typename FamilyType::MI_STORE_DATA_IMM; + using MI_SEMAPHORE_WAIT = typename FamilyType::MI_SEMAPHORE_WAIT; + using COMPUTE_WALKER = typename FamilyType::COMPUTE_WALKER; + using POSTSYNC_DATA = typename FamilyType::POSTSYNC_DATA; + + class MyMockCmdList : public WhiteBox> { + public: + using BaseClass = WhiteBox>; + using BaseClass::BaseClass; + + ze_result_t flushImmediate(ze_result_t inputRet, bool performMigration, bool hasStallingCmds, bool hasRelaxedOrderingDependencies, bool kernelOperation, ze_event_handle_t hSignalEvent) override { + flushData.push_back(this->cmdListCurrentStartOffset); + + this->cmdListCurrentStartOffset = this->commandContainer.getCommandStream()->getUsed(); + + return ZE_RESULT_SUCCESS; + } + + std::vector flushData; // start_offset + }; + + DebugManager.flags.DirectSubmissionRelaxedOrdering.set(1); + DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.set(1); + + auto ultCsr = static_cast *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver); + + auto directSubmission = new MockDirectSubmissionHw>(*ultCsr); + ultCsr->directSubmission.reset(directSubmission); + int client1, client2; + ultCsr->registerClient(&client1); + ultCsr->registerClient(&client2); + + auto immCmdList = createImmCmdListImpl(); + + auto cmdStream = immCmdList->getCmdContainer().getCommandStream(); + + auto eventPool = createEvents(1, true); + events[0]->signalScope = 0; + + immCmdList->inOrderDependencyCounter = 1; + + EXPECT_TRUE(immCmdList->isRelaxedOrderingDispatchAllowed(0)); + + EXPECT_EQ(0u, immCmdList->flushData.size()); + + zeCommandListAppendLaunchKernel(immCmdList->toHandle(), kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr); + + ASSERT_EQ(2u, immCmdList->flushData.size()); + EXPECT_EQ(2u, immCmdList->inOrderDependencyCounter); + + { + + GenCmdList cmdList; + ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, cmdStream->getCpuBase(), immCmdList->flushData[1])); + + auto sdiItor = find(cmdList.begin(), cmdList.end()); + ASSERT_NE(cmdList.end(), sdiItor); + + auto sdiCmd = genCmdCast(*sdiItor); + ASSERT_NE(nullptr, sdiCmd); + + EXPECT_EQ(events[0]->getCompletionFieldGpuAddress(device), sdiCmd->getAddress()); + EXPECT_EQ(0u, sdiCmd->getStoreQword()); + EXPECT_EQ(Event::STATE_CLEARED, sdiCmd->getDataDword0()); + + auto sdiOffset = ptrDiff(sdiCmd, cmdStream->getCpuBase()); + EXPECT_TRUE(sdiOffset >= immCmdList->flushData[0]); + EXPECT_TRUE(sdiOffset < immCmdList->flushData[1]); + + auto walkerItor = find(sdiItor, cmdList.end()); + ASSERT_NE(cmdList.end(), walkerItor); + + auto walkerCmd = genCmdCast(*walkerItor); + auto &postSync = walkerCmd->getPostSync(); + + auto eventBaseGpuVa = events[0]->getPacketAddress(device); + + EXPECT_EQ(POSTSYNC_DATA::OPERATION_WRITE_TIMESTAMP, postSync.getOperation()); + EXPECT_EQ(eventBaseGpuVa, postSync.getDestinationAddress()); + + auto walkerOffset = ptrDiff(walkerCmd, cmdStream->getCpuBase()); + EXPECT_TRUE(walkerOffset >= immCmdList->flushData[0]); + EXPECT_TRUE(walkerOffset < immCmdList->flushData[1]); + } + + { + + GenCmdList cmdList; + ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, ptrOffset(cmdStream->getCpuBase(), immCmdList->flushData[1]), (cmdStream->getUsed() - immCmdList->flushData[1]))); + + // Relaxed Ordering registers + auto lrrCmd = genCmdCast(*cmdList.begin()); + ASSERT_NE(nullptr, lrrCmd); + + EXPECT_EQ(CS_GPR_R4, lrrCmd->getSourceRegisterAddress()); + EXPECT_EQ(CS_GPR_R0, lrrCmd->getDestinationRegisterAddress()); + lrrCmd++; + EXPECT_EQ(CS_GPR_R4 + 4, lrrCmd->getSourceRegisterAddress()); + EXPECT_EQ(CS_GPR_R0 + 4, lrrCmd->getDestinationRegisterAddress()); + + lrrCmd++; + + auto eventEndGpuVa = events[0]->getCompletionFieldGpuAddress(device); + + EXPECT_TRUE(RelaxedOrderingCommandsHelper::verifyConditionalDataMemBbStart(lrrCmd, 0, eventEndGpuVa, static_cast(Event::STATE_CLEARED), + NEO::CompareOperation::Equal, true, false)); + + auto sdiCmd = genCmdCast(ptrOffset(lrrCmd, EncodeBatchBufferStartOrEnd::getCmdSizeConditionalDataMemBatchBufferStart(false))); + ASSERT_NE(nullptr, sdiCmd); + + EXPECT_EQ(immCmdList->inOrderDependencyCounterAllocation->getGpuAddress(), sdiCmd->getAddress()); + EXPECT_EQ(immCmdList->isQwordInOrderCounter(), sdiCmd->getStoreQword()); + EXPECT_EQ(2u, sdiCmd->getDataDword0()); + } +} + HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingRegularEventThenClearAndChainWithSyncAllocSignaling, IsAtLeastXeHpCore) { using MI_STORE_DATA_IMM = typename FamilyType::MI_STORE_DATA_IMM; using MI_SEMAPHORE_WAIT = typename FamilyType::MI_SEMAPHORE_WAIT; @@ -1661,6 +1785,19 @@ HWTEST2_F(InOrderCmdListTests, givenNonPostSyncWalkerWhenPatchingThenThrow, NonP EXPECT_ANY_THROW(walkerCmd.patch(1)); } +HWTEST2_F(InOrderCmdListTests, givenNonPostSyncWalkerWhenAskingForNonWalkerSignalingRequiredThenReturnFalse, NonPostSyncWalkerMatcher) { + auto immCmdList = createImmCmdList(); + + auto eventPool1 = createEvents(1, true); + auto eventPool2 = createEvents(1, false); + auto eventPool3 = createEvents(1, false); + events[2]->inOrderExecEvent = false; + + EXPECT_FALSE(immCmdList->isInOrderNonWalkerSignalingRequired(events[0].get())); + EXPECT_FALSE(immCmdList->isInOrderNonWalkerSignalingRequired(events[1].get())); + EXPECT_FALSE(immCmdList->isInOrderNonWalkerSignalingRequired(events[2].get())); +} + HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingWalkerThenProgramPipeControlWithSignalAllocation, NonPostSyncWalkerMatcher) { using WALKER = typename FamilyType::WALKER_TYPE; using PIPE_CONTROL = typename FamilyType::PIPE_CONTROL; diff --git a/shared/source/debug_settings/debug_variables_base.inl b/shared/source/debug_settings/debug_variables_base.inl index 4c4805cb40..d4291abced 100644 --- a/shared/source/debug_settings/debug_variables_base.inl +++ b/shared/source/debug_settings/debug_variables_base.inl @@ -254,6 +254,7 @@ DECLARE_DEBUG_VARIABLE(int32_t, WaitForUserFenceOnEventHostSynchronize, -1, "-1: DECLARE_DEBUG_VARIABLE(int32_t, DisableSystemPointerKernelArgument, -1, "-1: default, 0: Disabled, 1: using a system pointer for kernel argument returns an error.") DECLARE_DEBUG_VARIABLE(int32_t, ProgramUserInterruptOnResolvedDependency, -1, "-1: default, 0: Disabled, >=1: bitfield. 01b: program after semaphore, 10b: on signaling fence (non-walker append).") DECLARE_DEBUG_VARIABLE(int32_t, EnableInOrderRegularCmdListPatching, -1, "-1: default, 0: Disabled, 1: If set, patch counter value on execute call") +DECLARE_DEBUG_VARIABLE(int32_t, EnableInOrderRelaxedOrderingForEventsChaining, -1, "-1: default, 0: Disabled, 1: If set, send 2 immediate flushes to avoid stalling RelaxedOrdering Scheduler.") /*LOGGING FLAGS*/ DECLARE_DEBUG_VARIABLE(int32_t, PrintDriverDiagnostics, -1, "prints driver diagnostics messages to standard output, value corresponds to hint level") diff --git a/shared/test/common/test_files/igdrcl.config b/shared/test/common/test_files/igdrcl.config index 3812ef5848..97b1a245c4 100644 --- a/shared/test/common/test_files/igdrcl.config +++ b/shared/test/common/test_files/igdrcl.config @@ -550,4 +550,5 @@ DisableSystemPointerKernelArgument = -1 DoNotValidateDriverPath = 0 EnableInOrderRegularCmdListPatching = -1 ForceInOrderEvents = -1 +EnableInOrderRelaxedOrderingForEventsChaining = -1 # Please don't edit below this line