Media
/
compute-runtime
mirror of https://github.com/intel/compute-runtime.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Minor refactoring of CommandComputeKernel class 

- class cannot operate without kernel object
- improved ULTs

Change-Id: I4d1a6c3685f3908ce07154605aea649cae349d27
This commit is contained in:
hjnapiat 2018-05-11 14:03:03 +02:00 committed by sys_ocldev
parent 621a2dfcd1
commit eedde057ea
3 changed files with 47 additions and 23 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

13
runtime/helpers/task_information.cpp
View File

@ -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<Device> 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);

24
unit_tests/command_queue/enqueue_kernel_tests.cpp
View File

@ -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<FamilyType>(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<FamilyType>(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);
}

33
unit_tests/event/event_tests.cpp
View File

@ -462,13 +462,16 @@ TEST_F(InternalsEventTest, processBlockedCommandsKernelOperation) {
auto blockedCommandsData = new KernelOperation(std::unique_ptr<LinearStream>(cmdStream), UniqueIH(dsh),
UniqueIH(ioh), UniqueIH(ssh), *pCmdQ->getDevice().getMemoryManager());
MockKernelWithInternals mockKernelWithInternals(*pDevice);
auto pKernel = mockKernelWithInternals.mockKernel;
auto &csr = pDevice->getCommandStreamReceiver();
std::vector<Surface *> 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<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode);
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1);
event.setCommand(std::unique_ptr<Command>(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<LinearStream>(cmdStream), UniqueIH(dsh),
UniqueIH(ioh), UniqueIH(ssh), *pCmdQ->getDevice().getMemoryManager());
MockKernelWithInternals mockKernelWithInternals(*pDevice);
auto pKernel = mockKernelWithInternals.mockKernel;
auto &csr = pDevice->getCommandStreamReceiver();
std::vector<Surface *> v;
NullSurface *surface = new NullSurface;
v.push_back(surface);
PreemptionMode preemptionMode = pDevice->getPreemptionMode();
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode);
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1);
event.setCommand(std::unique_ptr<Command>(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<KernelInfo *>(&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<Surface *> v;
PreemptionMode preemptionMode = pDevice->getPreemptionMode();
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, std::move(printfHandler), preemptionMode, pKernel);
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, std::move(printfHandler), preemptionMode, pKernel, 1);
event.setCommand(std::unique_ptr<Command>(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<FamilyType>();
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<Surface *> v;
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode);
auto cmd = new CommandComputeKernel(*pCmdQ, csr, std::unique_ptr<KernelOperation>(blockedCommandsData), v, false, false, false, nullptr, preemptionMode, pKernel, 1);
event->setCommand(std::unique_ptr<Command>(cmd));
FlushStamp expectedFlushStamp = 0;