intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-12-28 00:03:14 +08:00
Code Issues Packages Projects Releases Wiki Activity

Move indirect heaps from command queues to csr.

Browse Source 
-This is required to enable N:1 submission model.
-If heaps are coming from different command queues that always
mean that STATE_BASE_ADDRESS needs to be reloaded
-In order to not emit any non pipelined state in CSR, this change
moves the ownership of IndirectHeap to one centralized place which is
CommandStreamReceiver
-This way when there are submissions from multiple command queues then
they reuse the same heaps, therefore preventing SBA reload

Change-Id: I5caf5dc5cb05d7a2d8766883d9bc51c29062e980
This commit is contained in:
Mrozek, Michal
2018-04-26 10:01:01 +02:00
committed by sys_ocldev
parent be7393fcfe
commit 8d2df3c332
13 changed files with 231 additions and 123 deletions
Download Patch File Download Diff File Expand all files Collapse all files

24
unit_tests/command_queue/command_queue_tests.cpp
View File

@@ -547,6 +547,7 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithoutHeapAllocationWhenA
MockCommandQueue cmdQ(context.get(), pDevice, props);
auto memoryManager = pDevice->getMemoryManager();
auto &csr = pDevice->getUltCommandStreamReceiver<DEFAULT_TEST_FAMILY_NAME>();
EXPECT_TRUE(memoryManager->allocationsForReuse.peekIsEmpty());
const auto &indirectHeap = cmdQ.getIndirectHeap(this->GetParam(), 100);
@@ -554,8 +555,8 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithoutHeapAllocationWhenA
auto graphicsAllocation = indirectHeap.getGraphicsAllocation();
cmdQ.indirectHeap[this->GetParam()]->replaceGraphicsAllocation(nullptr);
cmdQ.indirectHeap[this->GetParam()]->replaceBuffer(nullptr, 0);
csr.indirectHeap[this->GetParam()]->replaceGraphicsAllocation(nullptr);
csr.indirectHeap[this->GetParam()]->replaceBuffer(nullptr, 0);
// Request a larger heap than the first.
cmdQ.getIndirectHeap(this->GetParam(), heapSize + 6000);
@@ -564,17 +565,15 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithoutHeapAllocationWhenA
memoryManager->freeGraphicsMemory(graphicsAllocation);
}
TEST_P(CommandQueueIndirectHeapTest, givenCommandQueueWithResourceCachingActiveWhenQueueISDestroyedThenIndirectHeapIsPutOnReuseList) {
TEST_P(CommandQueueIndirectHeapTest, givenCommandQueueWithResourceCachingActiveWhenQueueISDestroyedThenIndirectHeapIsNotOnReuseList) {
auto cmdQ = new CommandQueue(context.get(), pDevice, 0);
auto memoryManager = pDevice->getMemoryManager();
const auto &indirectHeap = cmdQ->getIndirectHeap(this->GetParam(), 100);
auto graphicsAllocation = indirectHeap.getGraphicsAllocation();
cmdQ->getIndirectHeap(this->GetParam(), 100);
EXPECT_TRUE(memoryManager->allocationsForReuse.peekIsEmpty());
//now destroy command queue, heap should go to reusable list
delete cmdQ;
EXPECT_FALSE(memoryManager->allocationsForReuse.peekIsEmpty());
EXPECT_TRUE(memoryManager->allocationsForReuse.peekContains(*graphicsAllocation));
EXPECT_TRUE(memoryManager->allocationsForReuse.peekIsEmpty());
}
TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithHeapAllocatedWhenIndirectHeapIsReleasedThenHeapAllocationAndHeapBufferIsSetToNullptr) {
@@ -593,8 +592,9 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithHeapAllocatedWhenIndir
EXPECT_NE(nullptr, graphicsAllocation);
cmdQ.releaseIndirectHeap(this->GetParam());
auto &csr = pDevice->getUltCommandStreamReceiver<DEFAULT_TEST_FAMILY_NAME>();
EXPECT_EQ(nullptr, cmdQ.indirectHeap[this->GetParam()]->getGraphicsAllocation());
EXPECT_EQ(nullptr, csr.indirectHeap[this->GetParam()]->getGraphicsAllocation());
EXPECT_EQ(nullptr, indirectHeap.getCpuBase());
EXPECT_EQ(0u, indirectHeap.getMaxAvailableSpace());
@@ -605,8 +605,9 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithoutHeapAllocatedWhenIn
MockCommandQueue cmdQ(context.get(), pDevice, props);
cmdQ.releaseIndirectHeap(this->GetParam());
auto &csr = pDevice->getUltCommandStreamReceiver<DEFAULT_TEST_FAMILY_NAME>();
EXPECT_EQ(nullptr, cmdQ.indirectHeap[this->GetParam()]);
EXPECT_EQ(nullptr, csr.indirectHeap[this->GetParam()]);
}
TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithHeapWhenGraphicAllocationIsNullThenNothingOnReuseList) {
@@ -616,9 +617,10 @@ TEST_P(CommandQueueIndirectHeapTest, GivenCommandQueueWithHeapWhenGraphicAllocat
auto &ih = cmdQ.getIndirectHeap(this->GetParam(), 0u);
auto allocation = ih.getGraphicsAllocation();
EXPECT_NE(nullptr, allocation);
auto &csr = pDevice->getUltCommandStreamReceiver<DEFAULT_TEST_FAMILY_NAME>();
cmdQ.indirectHeap[this->GetParam()]->replaceGraphicsAllocation(nullptr);
cmdQ.indirectHeap[this->GetParam()]->replaceBuffer(nullptr, 0);
csr.indirectHeap[this->GetParam()]->replaceGraphicsAllocation(nullptr);
csr.indirectHeap[this->GetParam()]->replaceBuffer(nullptr, 0);
cmdQ.releaseIndirectHeap(this->GetParam());

19
unit_tests/command_queue/enqueue_kernel_tests.cpp
View File

@@ -1530,6 +1530,25 @@ HWTEST_F(EnqueueKernelTest, givenKernelWhenItIsEnqueuedThenAllResourceGraphicsAl
EXPECT_EQ(csrTaskCount, allocation->taskCount);
}
}
HWTEST_F(EnqueueKernelTest, givenKernelWhenItIsSubmittedFromTwoDifferentCommandQueuesThenCsrDoesntReloadAnyCommands) {
auto &csr = this->pDevice->getUltCommandStreamReceiver<FamilyType>();
MockKernelWithInternals mockKernel(*pDevice);
size_t gws[3] = {1, 0, 0};
pCmdQ->enqueueKernel(mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr);
auto currentUsed = csr.commandStream.getUsed();
const cl_queue_properties props[] = {0};
auto inOrderQueue = clCreateCommandQueueWithProperties(context, pDevice, props, nullptr);
clEnqueueNDRangeKernel(inOrderQueue, mockKernel.mockKernel, 1, nullptr, gws, nullptr, 0, nullptr, nullptr);
auto usedAfterSubmission = csr.commandStream.getUsed();
EXPECT_EQ(usedAfterSubmission, currentUsed);
clReleaseCommandQueue(inOrderQueue);
}
TEST_F(EnqueueKernelTest, givenKernelWhenAllArgsAreNotAndEventExistSetThenClEnqueueNDRangeKernelReturnsInvalidKernelArgsAndSetEventToNull) {
const size_t n = 512;
size_t globalWorkSize[3] = {n, 1, 1};

42
unit_tests/command_queue/oom_buffer_tests.cpp
View File

@@ -123,7 +123,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueCopyBuffer) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
}
HWTEST_P(OOMCommandQueueBufferTest, enqueueFillBuffer) {
@@ -143,7 +147,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueFillBuffer) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
}
HWTEST_P(OOMCommandQueueBufferTest, enqueueReadBuffer) {
@@ -163,7 +171,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueReadBuffer) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
}
HWTEST_P(OOMCommandQueueBufferTest, enqueueWriteBuffer) {
@@ -183,7 +195,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueWriteBuffer) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
}
HWTEST_P(OOMCommandQueueBufferTest, enqueueWriteBufferRect) {
@@ -203,7 +219,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueWriteBufferRect) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
}
HWTEST_P(OOMCommandQueueBufferTest, enqueueKernelHelloWorld) {
@@ -226,7 +246,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueKernelHelloWorld) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);
@@ -252,7 +276,11 @@ HWTEST_P(OOMCommandQueueBufferTest, enqueueKernelSimpleArg) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);

24
unit_tests/command_queue/oom_image_tests.cpp
View File

@@ -104,7 +104,11 @@ HWTEST_P(OOMCommandQueueImageTest, enqueueCopyImage) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);
@@ -124,7 +128,11 @@ HWTEST_P(OOMCommandQueueImageTest, enqueueFillImage) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);
@@ -144,7 +152,11 @@ HWTEST_P(OOMCommandQueueImageTest, enqueueReadImage) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);
@@ -164,7 +176,11 @@ HWTEST_P(OOMCommandQueueImageTest, enqueueWriteImage) {
auto usedAfterCS = commandStream.getUsed();
auto usedAfterISH = indirectHeap.getUsed();
EXPECT_LE(usedAfterCS - usedBeforeCS, commandStream.getMaxAvailableSpace());
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
if (usedAfterISH > usedBeforeISH) {
EXPECT_LE(usedAfterISH - usedBeforeISH, indirectHeap.getMaxAvailableSpace());
} else {
EXPECT_LE(usedAfterISH, indirectHeap.getMaxAvailableSpace());
}
EXPECT_EQ(CL_SUCCESS, retVal1);
EXPECT_EQ(CL_SUCCESS, retVal2);