diff --git a/runtime/command_queue/command_queue_hw.h b/runtime/command_queue/command_queue_hw.h index 258512c8c5..684564a3ad 100644 --- a/runtime/command_queue/command_queue_hw.h +++ b/runtime/command_queue/command_queue_hw.h @@ -66,6 +66,7 @@ class CommandQueueHw : public CommandQueue { if (getCmdQueueProperties(properties, CL_QUEUE_PROPERTIES) & static_cast(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)) { device->getCommandStreamReceiver().overrideDispatchPolicy(DispatchMode::BatchedDispatch); + device->getCommandStreamReceiver().enableNTo1SubmissionModel(); } } diff --git a/runtime/command_queue/enqueue_common.h b/runtime/command_queue/enqueue_common.h index fbacef1350..f7251964f9 100644 --- a/runtime/command_queue/enqueue_common.h +++ b/runtime/command_queue/enqueue_common.h @@ -556,7 +556,7 @@ CompletionStamp CommandQueueHw::enqueueNonBlocked( dispatchFlags.implicitFlush = implicitFlush; dispatchFlags.flushStampReference = this->flushStamp->getStampReference(); dispatchFlags.preemptionMode = PreemptionHelper::taskPreemptionMode(*device, multiDispatchInfo); - dispatchFlags.outOfOrderExecutionAllowed = !eventBuilder.getEvent() || this->isOOQEnabled(); + dispatchFlags.outOfOrderExecutionAllowed = !eventBuilder.getEvent() || commandStreamReceiver.isNTo1SubmissionModelEnabled(); DEBUG_BREAK_IF(taskLevel >= Event::eventNotReady); diff --git a/runtime/command_stream/command_stream_receiver.h b/runtime/command_stream/command_stream_receiver.h index 905e6169f2..0fd2a4218a 100644 --- a/runtime/command_stream/command_stream_receiver.h +++ b/runtime/command_stream/command_stream_receiver.h @@ -104,6 +104,8 @@ class CommandStreamReceiver { uint32_t peekLatestFlushedTaskCount() const { return latestFlushedTaskCount; } + void enableNTo1SubmissionModel() { this->nTo1SubmissionModelEnabled = true; } + bool isNTo1SubmissionModelEnabled() const { return this->nTo1SubmissionModelEnabled; } void overrideDispatchPolicy(DispatchMode overrideValue) { this->dispatchMode = overrideValue; } virtual void overrideMediaVFEStateDirty(bool dirty) { mediaVfeStateDirty = dirty; } @@ -173,6 +175,7 @@ class CommandStreamReceiver { std::unique_ptr osInterface; std::unique_ptr submissionAggregator; + bool nTo1SubmissionModelEnabled = false; DispatchMode dispatchMode = DispatchMode::ImmediateDispatch; bool disableL3Cache = false; uint32_t requiredScratchSize = 0; diff --git a/unit_tests/api/cl_create_command_queue_tests.cpp b/unit_tests/api/cl_create_command_queue_tests.cpp index 58b6d30ccd..12d7658ef7 100644 --- a/unit_tests/api/cl_create_command_queue_tests.cpp +++ b/unit_tests/api/cl_create_command_queue_tests.cpp @@ -84,4 +84,16 @@ HWTEST_F(clCreateCommandQueueTest, givenOoqParametersWhenQueueIsCreatedThenComma EXPECT_EQ(DispatchMode::BatchedDispatch, csr.dispatchMode); retVal = clReleaseCommandQueue(cmdq); } + +HWTEST_F(clCreateCommandQueueTest, givenOoqParametersWhenQueueIsCreatedThenCommandStreamReceiverSwitchesToNTo1SubmissionModel) { + cl_int retVal = CL_SUCCESS; + cl_queue_properties ooq = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE; + auto &csr = reinterpret_cast &>(pContext->getDevice(0)->getCommandStreamReceiver()); + EXPECT_FALSE(csr.isNTo1SubmissionModelEnabled()); + + auto cmdq = clCreateCommandQueue(pContext, devices[0], ooq, &retVal); + EXPECT_TRUE(csr.isNTo1SubmissionModelEnabled()); + retVal = clReleaseCommandQueue(cmdq); +} + } // namespace ULT diff --git a/unit_tests/command_queue/enqueue_kernel_tests.cpp b/unit_tests/command_queue/enqueue_kernel_tests.cpp index a40675515d..232502fa97 100644 --- a/unit_tests/command_queue/enqueue_kernel_tests.cpp +++ b/unit_tests/command_queue/enqueue_kernel_tests.cpp @@ -1469,10 +1469,34 @@ HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelReturningEv clReleaseEvent(event); } -HWTEST_F(EnqueueKernelTest, givenOutOfOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeThenPipeControlPositionIsRecorded) { - const cl_queue_properties props[3] = {CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, 0}; +HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeAndCommandStreamReceiverIsInNTo1ModeThenPipeControlPositionIsRecorded) { + const cl_queue_properties props[] = {0}; auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr); + auto mockCsr = new MockCsrHw2(pDevice->getHardwareInfo()); + mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch); + pDevice->resetCommandStreamReceiver(mockCsr); + mockCsr->enableNTo1SubmissionModel(); + + auto mockedSubmissionsAggregator = new mockSubmissionsAggregator(); + mockCsr->overrideSubmissionAggregator(mockedSubmissionsAggregator); + + MockKernelWithInternals mockKernel(*pDevice); + size_t gws[3] = {1, 0, 0}; + cl_event event; + + clEnqueueNDRangeKernel(inOrderQueue, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, &event); + + EXPECT_FALSE(mockedSubmissionsAggregator->peekCmdBufferList().peekIsEmpty()); + auto cmdBuffer = mockedSubmissionsAggregator->peekCmdBufferList().peekHead(); + EXPECT_NE(nullptr, cmdBuffer->pipeControlThatMayBeErasedLocation); + EXPECT_NE(nullptr, cmdBuffer->epiloguePipeControlLocation); + + clReleaseCommandQueue(inOrderQueue); + clReleaseEvent(event); +} + +HWTEST_F(EnqueueKernelTest, givenOutOfOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeThenPipeControlPositionIsRecorded) { auto mockCsr = new MockCsrHw2(pDevice->getHardwareInfo()); mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch); pDevice->resetCommandStreamReceiver(mockCsr); @@ -1480,6 +1504,9 @@ HWTEST_F(EnqueueKernelTest, givenOutOfOrderCommandQueueWhenEnqueueKernelReturnin auto mockedSubmissionsAggregator = new mockSubmissionsAggregator(); mockCsr->overrideSubmissionAggregator(mockedSubmissionsAggregator); + const cl_queue_properties props[3] = {CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, 0}; + auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr); + MockKernelWithInternals mockKernel(*pDevice); size_t gws[3] = {1, 0, 0}; cl_event event;