Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Enable out of order execution for all submissions.

Browse Source 
- This change enabled multiple independent command queues to execute
concurrently without stalling pipe controls in between
- This change removes L3 flushes between kernels
- Dependencies between commands are resolved via task level mechanism
- Out of order queues are not changing task level between submissions
- In order queues are increasing task level between submissions
- Whenever task level changes there is pipe control with cs stall emitted
between GPGPU_WALKERs

Change-Id: I558653b296424e4775d060df3072e2a50684b715
This commit is contained in:
Mrozek, Michal
2018-02-07 15:49:04 +01:00
committed by sys_ocldev
parent 043ad0eeba
commit d8f2142faa
2 changed files with 23 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
runtime/command_queue/enqueue_common.h
View File

@ -538,7 +538,7 @@ CompletionStamp CommandQueueHw<GfxFamily>::enqueueNonBlocked(
dispatchFlags.implicitFlush = implicitFlush;
dispatchFlags.flushStampReference = this->flushStamp->getStampReference();
dispatchFlags.preemptionMode = PreemptionHelper::taskPreemptionMode(*device, multiDispatchInfo);
dispatchFlags.outOfOrderExecutionAllowed = this->isOOQEnabled();
dispatchFlags.outOfOrderExecutionAllowed = true;
DEBUG_BREAK_IF(taskLevel >= Event::eventNotReady);

22
unit_tests/command_queue/enqueue_kernel_tests.cpp
View File

@ -1421,6 +1421,28 @@ HWTEST_F(EnqueueKernelTest, givenOutOfOrderCommandQueueWhenEnqueueKernelIsMadeTh
clReleaseCommandQueue(ooq);
}
HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelIsMadeThenPipeControlPositionIsRecorded) {
const cl_queue_properties props[] = {0};
auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr);
auto mockCsr = new MockCsrHw2<FamilyType>(pDevice->getHardwareInfo());
mockCsr->overrideDispatchPolicy(CommandStreamReceiver::DispatchMode::BatchedDispatch);
pDevice->resetCommandStreamReceiver(mockCsr);
auto mockedSubmissionsAggregator = new mockSubmissionsAggregator();
mockCsr->overrideSubmissionAggregator(mockedSubmissionsAggregator);
MockKernelWithInternals mockKernel(*pDevice);
size_t gws[3] = {1, 0, 0};
clEnqueueNDRangeKernel(inOrderQueue, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr);
EXPECT_FALSE(mockedSubmissionsAggregator->peekCmdBufferList().peekIsEmpty());
auto cmdBuffer = mockedSubmissionsAggregator->peekCmdBufferList().peekHead();
EXPECT_NE(nullptr, cmdBuffer->pipeControlLocation);
clReleaseCommandQueue(inOrderQueue);
}
HWTEST_F(EnqueueKernelTest, givenCsrInBatchingModeWhenBlockingCallIsMadeThenEventAssociatedWithCommandHasProperFlushStamp) {
DebugManagerStateRestore stateRestore;
DebugManager.flags.MakeEachEnqueueBlocking.set(true);