diff --git a/Jenkinsfile b/Jenkinsfile index 7e982a51ae..12d7af76c3 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1,5 +1,5 @@ #!groovy neoDependenciesRev='735095-769' strategy='EQUAL' -allowedF=44 -allowedCD=347 +allowedF=43 +allowedCD=345 diff --git a/runtime/command_queue/enqueue_common.h b/runtime/command_queue/enqueue_common.h index 000f26f794..157f146e39 100644 --- a/runtime/command_queue/enqueue_common.h +++ b/runtime/command_queue/enqueue_common.h @@ -538,7 +538,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() && !multiDispatchInfo.begin()->getKernel()->isUsingSharedObjArgs()) || this->isOOQEnabled(); DEBUG_BREAK_IF(taskLevel >= Event::eventNotReady); diff --git a/unit_tests/command_queue/enqueue_kernel_tests.cpp b/unit_tests/command_queue/enqueue_kernel_tests.cpp index 1f63a5250b..79d586a54f 100644 --- a/unit_tests/command_queue/enqueue_kernel_tests.cpp +++ b/unit_tests/command_queue/enqueue_kernel_tests.cpp @@ -1443,7 +1443,30 @@ HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelIsMadeThenP clReleaseCommandQueue(inOrderQueue); } -HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeThenPipeControlPositionIsRecorded) { +HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasSharedObjectsAsArgIsMadeThenPipeControlPositionIsNotRecorded) { + const cl_queue_properties props[] = {0}; + auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr); + + auto mockCsr = new MockCsrHw2(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}; + mockKernel.mockKernel->setUsingSharedArgs(true); + clEnqueueNDRangeKernel(inOrderQueue, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr); + + EXPECT_FALSE(mockedSubmissionsAggregator->peekCmdBufferList().peekIsEmpty()); + auto cmdBuffer = mockedSubmissionsAggregator->peekCmdBufferList().peekHead(); + EXPECT_EQ(nullptr, cmdBuffer->pipeControlLocation); + + clReleaseCommandQueue(inOrderQueue); +} + +HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeThenPipeControlPositionIsNotRecorded) { const cl_queue_properties props[] = {0}; auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr); diff --git a/unit_tests/mocks/mock_kernel.h b/unit_tests/mocks/mock_kernel.h index cd0cb38745..709b3639af 100644 --- a/unit_tests/mocks/mock_kernel.h +++ b/unit_tests/mocks/mock_kernel.h @@ -180,6 +180,8 @@ class MockKernel : public Kernel { // Make protected members from base class publicly accessible in mock class using Kernel::kernelArgHandlers; + + void setUsingSharedArgs(bool usingSharedArgValue) { this->usingSharedObjArgs = usingSharedArgValue; } }; //class below have enough internals to service Enqueue operation.