From eedde057ea282787526733af9e69f7fdaf249a08 Mon Sep 17 00:00:00 2001 From: hjnapiat Date: Fri, 11 May 2018 14:03:03 +0200 Subject: [PATCH] Minor refactoring of CommandComputeKernel class - class cannot operate without kernel object - improved ULTs Change-Id: I4d1a6c3685f3908ce07154605aea649cae349d27 --- runtime/helpers/task_information.cpp | 13 +++----- .../command_queue/enqueue_kernel_tests.cpp | 24 ++++++++++++++ unit_tests/event/event_tests.cpp | 33 ++++++++++--------- 3 files changed, 47 insertions(+), 23 deletions(-) diff --git a/runtime/helpers/task_information.cpp b/runtime/helpers/task_information.cpp index e53e6e410c..0db2a7a983 100644 --- a/runtime/helpers/task_information.cpp +++ b/runtime/helpers/task_information.cpp @@ -121,9 +121,8 @@ CommandComputeKernel::CommandComputeKernel(CommandQueue &commandQueue, CommandSt this->surfaces.push_back(surface); } this->kernel = kernel; - if (kernel) { - kernel->incRefInternal(); - } + UNRECOVERABLE_IF(nullptr == this->kernel); + kernel->incRefInternal(); this->kernelCount = kernelCount; this->preemptionMode = preemptionMode; } @@ -136,16 +135,14 @@ CommandComputeKernel::~CommandComputeKernel() { if (kernelOperation->ioh.get() == kernelOperation->dsh.get()) { kernelOperation->doNotFreeISH = true; } - if (kernel) { - kernel->decRefInternal(); - } + kernel->decRefInternal(); } CompletionStamp &CommandComputeKernel::submit(uint32_t taskLevel, bool terminated) { if (terminated) { return completionStamp; } - bool executionModelKernel = kernel != nullptr ? kernel->isParentKernel : false; + bool executionModelKernel = kernel->isParentKernel; auto devQueue = commandQueue.getContext().getDefaultDeviceQueue(); TakeOwnershipWrapper deviceOwnership(commandQueue.getDevice()); @@ -223,7 +220,7 @@ CompletionStamp &CommandComputeKernel::submit(uint32_t taskLevel, bool terminate dispatchFlags.lowPriority = commandQueue.getPriority() == QueuePriority::LOW; dispatchFlags.throttle = commandQueue.getThrottle(); dispatchFlags.preemptionMode = preemptionMode; - dispatchFlags.mediaSamplerRequired = (kernel != nullptr) ? kernel->isVmeKernel() : false; + dispatchFlags.mediaSamplerRequired = kernel->isVmeKernel(); 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 89809b115a..0a08b322e4 100644 --- a/unit_tests/command_queue/enqueue_kernel_tests.cpp +++ b/unit_tests/command_queue/enqueue_kernel_tests.cpp @@ -1587,3 +1587,27 @@ TEST_F(EnqueueKernelTest, givenEnqueueCommandThatLwsExceedsDeviceCapabilitiesWhe auto status = pCmdQ->enqueueKernel(mockKernel.mockKernel, 1, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr); EXPECT_EQ(CL_INVALID_WORK_GROUP_SIZE, status); } + +HWTEST_F(EnqueueKernelTest, givenVMEKernelWhenEnqueueKernelThenDispatchFlagsHaveMediaSamplerRequired) { + auto mockCsr = new MockCsrHw2(pDevice->getHardwareInfo()); + mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch); + pDevice->resetCommandStreamReceiver(mockCsr); + + MockKernelWithInternals mockKernel(*pDevice, context); + size_t gws[3] = {1, 0, 0}; + mockKernel.kernelInfo.isVmeWorkload = true; + clEnqueueNDRangeKernel(this->pCmdQ, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr); + EXPECT_TRUE(mockCsr->passedDispatchFlags.mediaSamplerRequired); +} + +HWTEST_F(EnqueueKernelTest, givenNonVMEKernelWhenEnqueueKernelThenDispatchFlagsDoesntHaveMediaSamplerRequired) { + auto mockCsr = new MockCsrHw2(pDevice->getHardwareInfo()); + mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch); + pDevice->resetCommandStreamReceiver(mockCsr); + + MockKernelWithInternals mockKernel(*pDevice, context); + size_t gws[3] = {1, 0, 0}; + mockKernel.kernelInfo.isVmeWorkload = false; + clEnqueueNDRangeKernel(this->pCmdQ, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr); + EXPECT_FALSE(mockCsr->passedDispatchFlags.mediaSamplerRequired); +} diff --git a/unit_tests/event/event_tests.cpp b/unit_tests/event/event_tests.cpp index 18b55174c6..00753e921c 100644 --- a/unit_tests/event/event_tests.cpp +++ b/unit_tests/event/event_tests.cpp @@ -462,13 +462,16 @@ TEST_F(InternalsEventTest, processBlockedCommandsKernelOperation) { auto blockedCommandsData = new KernelOperation(std::unique_ptr(cmdStream), UniqueIH(dsh), UniqueIH(ioh), UniqueIH(ssh), *pCmdQ->getDevice().getMemoryManager()); + MockKernelWithInternals mockKernelWithInternals(*pDevice); + auto pKernel = mockKernelWithInternals.mockKernel; + auto &csr = pDevice->getCommandStreamReceiver(); std::vector v; SurfaceMock *surface = new SurfaceMock; surface->graphicsAllocation = new GraphicsAllocation((void *)0x1234, 100u); PreemptionMode preemptionMode = pDevice->getPreemptionMode(); v.push_back(surface); - auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode); + auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1); event.setCommand(std::unique_ptr(cmd)); auto taskLevelBefore = csr.peekTaskLevel(); @@ -478,7 +481,6 @@ TEST_F(InternalsEventTest, processBlockedCommandsKernelOperation) { auto taskLevelAfter = csr.peekTaskLevel(); EXPECT_EQ(taskLevelBefore + 1, taskLevelAfter); - delete pCmdQ; EXPECT_EQ(surface->resident, 1u); @@ -500,12 +502,15 @@ TEST_F(InternalsEventTest, processBlockedCommandsAbortKernelOperation) { auto blockedCommandsData = new KernelOperation(std::unique_ptr(cmdStream), UniqueIH(dsh), UniqueIH(ioh), UniqueIH(ssh), *pCmdQ->getDevice().getMemoryManager()); + MockKernelWithInternals mockKernelWithInternals(*pDevice); + auto pKernel = mockKernelWithInternals.mockKernel; + auto &csr = pDevice->getCommandStreamReceiver(); std::vector v; NullSurface *surface = new NullSurface; v.push_back(surface); PreemptionMode preemptionMode = pDevice->getPreemptionMode(); - auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode); + auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1); event.setCommand(std::unique_ptr(cmd)); auto taskLevelBefore = csr.peekTaskLevel(); @@ -544,14 +549,12 @@ TEST_F(InternalsEventTest, givenBlockedKernelWithPrintfWhenSubmittedThenPrintOut printfStringInfo.SizeInBytes = sizeof("test"); printfStringInfo.pStringData = testString; - KernelInfo *pKernelInfo = new KernelInfo(); - pKernelInfo->patchInfo.pAllocateStatelessPrintfSurface = pPrintfSurface; - pKernelInfo->patchInfo.stringDataMap.insert(std::make_pair(0, printfStringInfo)); - - MockProgram *pProgram = new MockProgram(mockContext, false); - + MockKernelWithInternals mockKernelWithInternals(*pDevice); + auto pKernel = mockKernelWithInternals.mockKernel; + KernelInfo *kernelInfo = const_cast(&pKernel->getKernelInfo()); + kernelInfo->patchInfo.pAllocateStatelessPrintfSurface = pPrintfSurface; + kernelInfo->patchInfo.stringDataMap.insert(std::make_pair(0, printfStringInfo)); uint64_t crossThread[10]; - MockKernel *pKernel = new MockKernel(pProgram, *pKernelInfo, *pDevice); pKernel->setCrossThreadData(&crossThread, sizeof(uint64_t) * 8); MockMultiDispatchInfo multiDispatchInfo(pKernel); @@ -566,7 +569,7 @@ TEST_F(InternalsEventTest, givenBlockedKernelWithPrintfWhenSubmittedThenPrintOut auto &csr = pDevice->getCommandStreamReceiver(); std::vector v; PreemptionMode preemptionMode = pDevice->getPreemptionMode(); - auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, std::move(printfHandler), preemptionMode, pKernel); + auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, std::move(printfHandler), preemptionMode, pKernel, 1); event.setCommand(std::unique_ptr(cmd)); event.submitCommand(false); @@ -576,9 +579,6 @@ TEST_F(InternalsEventTest, givenBlockedKernelWithPrintfWhenSubmittedThenPrintOut EXPECT_FALSE(surface->isResident()); delete pPrintfSurface; - delete pKernelInfo; - pKernel->decRefInternal(); - pProgram->decRefInternal(); delete pCmdQ; } @@ -1436,6 +1436,9 @@ HWTEST_F(InternalsEventTest, givenAbortedCommandWhenSubmitCalledThenDontUpdateFl auto &csr = pDevice->getUltCommandStreamReceiver(); csr.flushStamp->setStamp(5); + MockKernelWithInternals mockKernelWithInternals(*pDevice); + auto pKernel = mockKernelWithInternals.mockKernel; + auto cmdStream = new LinearStream(alignedMalloc(4096, 4096), 4096); IndirectHeap *dsh = nullptr, *ioh = nullptr, *ssh = nullptr; pCmdQ->allocateHeapMemory(IndirectHeap::DYNAMIC_STATE, 4096u, dsh); @@ -1446,7 +1449,7 @@ HWTEST_F(InternalsEventTest, givenAbortedCommandWhenSubmitCalledThenDontUpdateFl UniqueIH(ioh), UniqueIH(ssh), *pCmdQ->getDevice().getMemoryManager()); PreemptionMode preemptionMode = pDevice->getPreemptionMode(); std::vector v; - auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode); + auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1); event->setCommand(std::unique_ptr(cmd)); FlushStamp expectedFlushStamp = 0;