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

feature: initial support for RelaxedOrdering of in-order Events chaining

Browse Source 
Disabled by default.

Related-To: NEO-7966

Signed-off-by: Dunajski, Bartosz <bartosz.dunajski@intel.com>
This commit is contained in:
Dunajski, Bartosz
2023-09-26 12:23:14 +00:00
committed by Compute-Runtime-Automation
parent 5dc56c221f
commit 4e8600d8d0
12 changed files with 207 additions and 14 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
level_zero/api/core/ze_module_api_entrypoints.h
View File

@@ -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(

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

@@ -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);

3
level_zero/core/source/cmdlist/cmdlist_hw.h
View File

@@ -68,6 +68,8 @@ template <GFXCORE_FAMILY gfxCoreFamily>
struct CommandListCoreFamily : CommandListImp {
using GfxFamily = typename NEO::GfxFamilyMapper<gfxCoreFamily>::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);

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

@@ -354,7 +354,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(ze_kernel_h
auto res = appendLaunchKernelWithParams(Kernel::fromHandle(kernelHandle), threadGroupDimensions,
event, launchParams);
if (isInOrderExecutionEnabled()) {
if (isInOrderExecutionEnabled() && !launchParams.skipInOrderNonWalkerSignaling) {
handleInOrderDependencyCounter();
}

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

@@ -165,7 +165,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
void checkAvailableSpace(uint32_t numEvents, bool hasRelaxedOrderingDependencies, size_t commandSize);
void updateDispatchFlagsWithRequiredStreamState(NEO::DispatchFlags &dispatchFlags);
ze_result_t flushImmediate(ze_result_t inputRet, bool performMigration, bool hasStallingCmds, bool hasRelaxedOrderingDependencies, bool kernelOperation, ze_event_handle_t hSignalEvent);
MOCKABLE_VIRTUAL ze_result_t flushImmediate(ze_result_t inputRet, bool performMigration, bool hasStallingCmds, bool hasRelaxedOrderingDependencies, bool kernelOperation, ze_event_handle_t hSignalEvent);
bool preferCopyThroughLockedPtr(CpuMemCopyInfo &cpuMemCopyInfo, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents);
bool isSuitableUSMHostAlloc(NEO::SvmAllocationData *alloc);
@@ -179,6 +179,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
size_t getTransferThreshold(TransferType transferType);
bool isBarrierRequired();
bool isRelaxedOrderingDispatchAllowed(uint32_t numWaitEvents) const override;
bool skipInOrderNonWalkerSignalingAllowed(ze_event_handle_t signalEvent) const override;
protected:
using BaseClass::inOrderDependencyCounter;
@@ -191,6 +192,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
void setupFlushMethod(const NEO::RootDeviceEnvironment &rootDeviceEnvironment) override;
bool isSkippingInOrderBarrierAllowed(ze_event_handle_t hSignalEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents) const;
void allocateOrReuseKernelPrivateMemoryIfNeeded(Kernel *kernel, uint32_t sizePerHwThread) override;
void handleInOrderNonWalkerSignaling(Event *event, bool &hasStallingCmds, bool &relaxedOrderingDispatch, ze_result_t &result);
MOCKABLE_VIRTUAL void checkAssert();
ComputeFlushMethodType computeFlushMethod = nullptr;

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

@@ -407,6 +407,11 @@ bool CommandListCoreFamilyImmediate<gfxCoreFamily>::hasStallingCmdsForRelaxedOrd
return (!relaxedOrderingDispatch && (numWaitEvents > 0 || this->inOrderDependencyCounter > 0));
}
template <GFXCORE_FAMILY gfxCoreFamily>
bool CommandListCoreFamilyImmediate<gfxCoreFamily>::skipInOrderNonWalkerSignalingAllowed(ze_event_handle_t signalEvent) const {
return this->isInOrderNonWalkerSignalingRequired(Event::fromHandle(signalEvent));
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchKernel(
ze_kernel_handle_t kernelHandle, const ze_group_count_t *threadGroupDimensions,
@@ -414,6 +419,7 @@ ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::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<gfxCoreFamily>::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<gfxCoreFamily>::handleInOrderDependencyCounter();
}
return flushImmediate(ret, true, stallingCmdsForRelaxedOrdering, relaxedOrderingDispatch, true, hSignalEvent);
}
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) {
nonWalkerSignalingHasRelaxedOrdering = isRelaxedOrderingDispatchAllowed(1);
}
if (nonWalkerSignalingHasRelaxedOrdering) {
result = flushImmediate(result, true, hasStallingCmds, relaxedOrderingDispatch, true, nullptr);
NEO::RelaxedOrderingHelper::encodeRegistersBeforeDependencyCheckers<GfxFamily>(*this->commandContainer.getCommandStream());
relaxedOrderingDispatch = true;
hasStallingCmds = hasStallingCmdsForRelaxedOrdering(1, relaxedOrderingDispatch);
}
CommandListCoreFamily<gfxCoreFamily>::appendWaitOnSingleEvent(event, nonWalkerSignalingHasRelaxedOrdering);
CommandListCoreFamily<gfxCoreFamily>::appendSignalInOrderDependencyCounter();
}
template <GFXCORE_FAMILY gfxCoreFamily>

6
level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl
View File

@@ -40,12 +40,18 @@ size_t CommandListCoreFamily<gfxCoreFamily>::getReserveSshSize() {
template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamily<gfxCoreFamily>::adjustWriteKernelTimestamp(uint64_t globalAddress, uint64_t contextAddress, bool maskLsb, uint32_t mask, bool workloadPartition) {}
template <GFXCORE_FAMILY gfxCoreFamily>
bool CommandListCoreFamily<gfxCoreFamily>::isInOrderNonWalkerSignalingRequired(const Event *event) const {
return false;
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::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<DriverHandleImp *>(device->getDriverHandle());
const auto &kernelDescriptor = kernel->getKernelDescriptor();
if (kernelDescriptor.kernelAttributes.flags.isInvalid) {

19
level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl
View File

@@ -77,9 +77,12 @@ void programEventL3Flush(Event *event,
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
Event *event,
bool CommandListCoreFamily<gfxCoreFamily>::isInOrderNonWalkerSignalingRequired(const Event *event) const {
return (event && (event->isUsingContextEndOffset() || !event->isInOrderExecEvent()));
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::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<gfxCoreFamily>::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<gfxCoreFamily>::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<gfxCoreFamily>::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<gfxCoreFamily>::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);

1
level_zero/core/test/unit_tests/mocks/mock_cmdlist.h
View File

@@ -176,6 +176,7 @@ struct WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>>
using BaseClass::inOrderPatchCmds;
using BaseClass::isBcsSplitNeeded;
using BaseClass::isFlushTaskSubmissionEnabled;
using BaseClass::isInOrderNonWalkerSignalingRequired;
using BaseClass::isQwordInOrderCounter;
using BaseClass::isSyncModeQueue;
using BaseClass::isTbxMode;

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

@@ -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 <GFXCORE_FAMILY gfxCoreFamily>
DestroyableZeUniquePtr<WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>>> createImmCmdList() {
auto cmdList = makeZeUniquePtr<WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>>>();
return createImmCmdListImpl<gfxCoreFamily, WhiteBox<L0::CommandListCoreFamilyImmediate<gfxCoreFamily>>>();
}
template <GFXCORE_FAMILY gfxCoreFamily, typename CmdListT>
DestroyableZeUniquePtr<CmdListT> createImmCmdListImpl() {
auto cmdList = makeZeUniquePtr<CmdListT>();
auto csr = device->getNEODevice()->getDefaultEngine().commandStreamReceiver;
@@ -1517,7 +1524,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingTimestampEventThen
auto eventPool = createEvents<FamilyType>(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<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 {
flushData.push_back(this->cmdListCurrentStartOffset);
this->cmdListCurrentStartOffset = this->commandContainer.getCommandStream()->getUsed();
return ZE_RESULT_SUCCESS;
}
std::vector<size_t> flushData; // start_offset
};
DebugManager.flags.DirectSubmissionRelaxedOrdering.set(1);
DebugManager.flags.EnableInOrderRelaxedOrderingForEventsChaining.set(1);
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 cmdStream = immCmdList->getCmdContainer().getCommandStream();
auto eventPool = createEvents<FamilyType>(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<MI_STORE_DATA_IMM *>(cmdList.begin(), cmdList.end());
ASSERT_NE(cmdList.end(), sdiItor);
auto sdiCmd = genCmdCast<MI_STORE_DATA_IMM *>(*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<COMPUTE_WALKER *>(sdiItor, cmdList.end());
ASSERT_NE(cmdList.end(), walkerItor);
auto walkerCmd = genCmdCast<COMPUTE_WALKER *>(*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<typename FamilyType::MI_LOAD_REGISTER_REG *>(*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<FamilyType>(lrrCmd, 0, eventEndGpuVa, static_cast<uint64_t>(Event::STATE_CLEARED),
NEO::CompareOperation::Equal, true, false));
auto sdiCmd = genCmdCast<MI_STORE_DATA_IMM *>(ptrOffset(lrrCmd, EncodeBatchBufferStartOrEnd<FamilyType>::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<gfxCoreFamily>();
auto eventPool1 = createEvents<FamilyType>(1, true);
auto eventPool2 = createEvents<FamilyType>(1, false);
auto eventPool3 = createEvents<FamilyType>(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;