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: enable RelaxedOrdering for in-order Events chaining

Browse Source 
Related-To: NEO-7966

Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
This commit is contained in:
Dunajski, Bartosz
2023-09-27 16:50:26 +00:00
committed by Compute-Runtime-Automation
parent 5f846d8a13
commit 2aaf5a1f03
2 changed files with 44 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -448,7 +448,7 @@ template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamilyImmediate<gfxCoreFamily>::handleInOrderNonWalkerSignaling(Event *event, bool &hasStallingCmds, bool &relaxedOrderingDispatch, ze_result_t &result) {
bool nonWalkerSignalingHasRelaxedOrdering = false;
if (NEO::DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.get() == 1) {
if (NEO::DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.get() != 0) {
nonWalkerSignalingHasRelaxedOrdering = isRelaxedOrderingDispatchAllowed(1);
}

44
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp
View File

@@ -1589,7 +1589,6 @@ HWTEST2_F(InOrderCmdListTests, givenRelaxedOrderingWhenProgrammingTimestampEvent
};
DebugManager.flags.DirectSubmissionRelaxedOrdering.set(1);
DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.set(1);
auto ultCsr = static_cast<UltCommandStreamReceiver<FamilyType> *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
@@ -1683,6 +1682,49 @@ HWTEST2_F(InOrderCmdListTests, givenRelaxedOrderingWhenProgrammingTimestampEvent
}
}
HWTEST2_F(InOrderCmdListTests, givenDebugFlagSetWhenChainingWithRelaxedOrderingThenSignalAsSingleSubmission, IsAtLeastXeHpcCore) {
class MyMockCmdList : public WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>> {
public:
using BaseClass = WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>>;
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 {
flushCount++;
return ZE_RESULT_SUCCESS;
}
uint32_t flushCount = 0;
};
DebugManager.flags.DirectSubmissionRelaxedOrdering.set(1);
DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.set(0);
auto ultCsr = static_cast<UltCommandStreamReceiver<FamilyType> *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
auto directSubmission = new MockDirectSubmissionHw<FamilyType, RenderDispatcher<FamilyType>>(*ultCsr);
ultCsr->directSubmission.reset(directSubmission);
int client1, client2;
ultCsr->registerClient(&client1);
ultCsr->registerClient(&client2);
auto immCmdList = createImmCmdListImpl<gfxCoreFamily, MyMockCmdList>();
auto eventPool = createEvents<FamilyType>(1, true);
events[0]->signalScope = 0;
immCmdList->inOrderDependencyCounter = 1;
EXPECT_TRUE(immCmdList->isRelaxedOrderingDispatchAllowed(0));
EXPECT_EQ(0u, immCmdList->flushCount);
zeCommandListAppendLaunchKernel(immCmdList->toHandle(), kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr);
ASSERT_EQ(1u, immCmdList->flushCount);
EXPECT_EQ(2u, immCmdList->inOrderDependencyCounter);
}
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;