From 1fce275542235ef72b719fca99bca98a0faa7fda Mon Sep 17 00:00:00 2001 From: "Dunajski, Bartosz" Date: Fri, 2 Mar 2018 16:32:14 +0100 Subject: [PATCH] Remove forced DC flush and disabled out of order execution for shared objects Change-Id: I0de86c3d5af488a347e83858f5dddbac2ef53c17 --- runtime/command_queue/enqueue_common.h | 4 ++-- unit_tests/command_queue/enqueue_kernel_tests.cpp | 8 ++++---- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/runtime/command_queue/enqueue_common.h b/runtime/command_queue/enqueue_common.h index b97aac4027..544a65bc45 100644 --- a/runtime/command_queue/enqueue_common.h +++ b/runtime/command_queue/enqueue_common.h @@ -539,7 +539,7 @@ CompletionStamp CommandQueueHw::enqueueNonBlocked( DispatchFlags dispatchFlags; dispatchFlags.blocking = blocking; - dispatchFlags.dcFlush = shouldFlushDC(commandType, printfHandler) || multiDispatchInfo.begin()->getKernel()->isUsingSharedObjArgs(); + dispatchFlags.dcFlush = shouldFlushDC(commandType, printfHandler); dispatchFlags.useSLM = slmUsed; dispatchFlags.guardCommandBufferWithPipeControl = true; dispatchFlags.GSBA32BitRequired = commandType == CL_COMMAND_NDRANGE_KERNEL; @@ -549,7 +549,7 @@ CompletionStamp CommandQueueHw::enqueueNonBlocked( dispatchFlags.implicitFlush = implicitFlush; dispatchFlags.flushStampReference = this->flushStamp->getStampReference(); dispatchFlags.preemptionMode = PreemptionHelper::taskPreemptionMode(*device, multiDispatchInfo); - dispatchFlags.outOfOrderExecutionAllowed = (!eventBuilder.getEvent() && !multiDispatchInfo.begin()->getKernel()->isUsingSharedObjArgs()) || this->isOOQEnabled(); + dispatchFlags.outOfOrderExecutionAllowed = !eventBuilder.getEvent() || 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 d140e0c76f..f521352c1a 100644 --- a/unit_tests/command_queue/enqueue_kernel_tests.cpp +++ b/unit_tests/command_queue/enqueue_kernel_tests.cpp @@ -1443,7 +1443,7 @@ HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelIsMadeThenP clReleaseCommandQueue(inOrderQueue); } -HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasSharedObjectsAsArgIsMadeThenPipeControlPositionIsNotRecorded) { +HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasSharedObjectsAsArgIsMadeThenPipeControlPositionIsRecorded) { const cl_queue_properties props[] = {0}; auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr); @@ -1461,13 +1461,13 @@ HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasShar EXPECT_FALSE(mockedSubmissionsAggregator->peekCmdBufferList().peekIsEmpty()); auto cmdBuffer = mockedSubmissionsAggregator->peekCmdBufferList().peekHead(); - EXPECT_EQ(nullptr, cmdBuffer->pipeControlThatMayBeErasedLocation); + EXPECT_NE(nullptr, cmdBuffer->pipeControlThatMayBeErasedLocation); EXPECT_NE(nullptr, cmdBuffer->epiloguePipeControlLocation); clReleaseCommandQueue(inOrderQueue); } -HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasSharedObjectsAsArgIsMadeThenPipeControlHasDcFlushSet) { +HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasSharedObjectsAsArgIsMadeThenPipeControlDoesntHaveDcFlush) { auto mockCsr = new MockCsrHw2(pDevice->getHardwareInfo()); mockCsr->overrideDispatchPolicy(CommandStreamReceiver::DispatchMode::BatchedDispatch); pDevice->resetCommandStreamReceiver(mockCsr); @@ -1477,7 +1477,7 @@ HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelThatHasShar mockKernel.mockKernel->setUsingSharedArgs(true); clEnqueueNDRangeKernel(this->pCmdQ, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr); - EXPECT_TRUE(mockCsr->passedDispatchFlags.dcFlush); + EXPECT_FALSE(mockCsr->passedDispatchFlags.dcFlush); } HWTEST_F(EnqueueKernelTest, givenInOrderCommandQueueWhenEnqueueKernelReturningEventIsMadeThenPipeControlPositionIsNotRecorded) {