Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-15 13:01:45 +08:00
Code Issues Packages Projects Releases Wiki Activity

Disable GTPin notifications for parent kernels

Browse Source 
Resolves: NEO-4877

Change-Id: I34d036d2c4db6976297456b889bb5dbbbd6182e2
Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>
This commit is contained in:
Mateusz Hoppe
2020-08-27 16:11:44 +02:00
committed by sys_ocldev
parent 8f52561307
commit a7e15b250c
3 changed files with 57 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

4
opencl/source/gtpin/gtpin_callbacks.cpp
View File

@ -68,7 +68,7 @@ void gtpinNotifyKernelCreate(cl_kernel kernel) {
auto &device = pKernel->getDevice();
GFXCORE_FAMILY genFamily = device.getHardwareInfo().platform.eRenderCoreFamily;
GTPinHwHelper &gtpinHelper = GTPinHwHelper::get(genFamily);
if (!gtpinHelper.addSurfaceState(pKernel)) {
if (pKernel->isParentKernel || !gtpinHelper.addSurfaceState(pKernel)) {
// Kernel with no SSH or Kernel EM, not supported
return;
}
@ -104,7 +104,7 @@ void gtpinNotifyKernelCreate(cl_kernel kernel) {
void gtpinNotifyKernelSubmit(cl_kernel kernel, void *pCmdQueue) {
if (isGTPinInitialized) {
auto pKernel = castToObjectOrAbort<Kernel>(kernel);
if (pKernel->getSurfaceStateHeapSize() == 0) {
if (pKernel->isParentKernel || pKernel->getSurfaceStateHeapSize() == 0) {
// Kernel with no SSH, not supported
return;
}

52
opencl/test/unit_test/gtpin/gtpin_tests.cpp
View File

@ -34,9 +34,11 @@
#include "opencl/test/unit_test/mocks/mock_buffer.h"
#include "opencl/test/unit_test/mocks/mock_command_queue.h"
#include "opencl/test/unit_test/mocks/mock_context.h"
#include "opencl/test/unit_test/mocks/mock_device_queue.h"
#include "opencl/test/unit_test/mocks/mock_kernel.h"
#include "opencl/test/unit_test/mocks/mock_platform.h"
#include "opencl/test/unit_test/program/program_tests.h"
#include "opencl/test/unit_test/test_macros/test_checks_ocl.h"
#include "test.h"
#include "gtest/gtest.h"
@ -1037,7 +1039,9 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithoutSSHIsUsedThenK
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithExecEnvIsUsedThenKernelCreateCallbacksIsNotCalled) {
HWCMDTEST_F(IGFX_GEN8_CORE, GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithDeviceEnqueueIsUsedThenKernelCreateAndSubmitCallbacksAreNotCalled) {
REQUIRE_DEVICE_ENQUEUE_OR_SKIP(pDevice);
gtpinCallbacks.onContextCreate = OnContextCreate;
gtpinCallbacks.onContextDestroy = OnContextDestroy;
gtpinCallbacks.onKernelCreate = OnKernelCreate;
@ -1053,6 +1057,17 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithExecEnvIsUsedThen
EXPECT_NE(nullptr, context);
auto pContext = castToObject<Context>(context);
cl_queue_properties devQproperties = 0;
auto devQ = std::make_unique<DeviceQueueHw<FamilyType>>(pContext, pDevice, devQproperties);
pContext->setDefaultDeviceQueue(devQ.get());
cl_command_queue cmdQ = nullptr;
cl_queue_properties properties = 0;
cmdQ = clCreateCommandQueue(context, device, properties, &retVal);
ASSERT_NE(nullptr, cmdQ);
EXPECT_EQ(CL_SUCCESS, retVal);
// Prepare a kernel with fake Execution Environment
char binary[1024] = {1, 2, 3, 4, 5, 6, 7, 8, 9, '\0'};
size_t binSize = 10;
@ -1153,6 +1168,20 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithExecEnvIsUsedThen
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(prevCount, KernelCreateCallbackCount);
int prevCount2 = KernelSubmitCallbackCount;
cl_uint workDim = 1;
size_t globalWorkOffset[3] = {0, 0, 0};
size_t globalWorkSize[3] = {1, 1, 1};
size_t localWorkSize[3] = {1, 1, 1};
MockParentKernel *parentKernel = MockParentKernel::create(*pContext);
retVal = clEnqueueNDRangeKernel(cmdQ, parentKernel, workDim, globalWorkOffset, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, retVal);
EXPECT_EQ(prevCount2, KernelSubmitCallbackCount);
delete parentKernel;
// Cleanup
retVal = clReleaseKernel(kernel);
EXPECT_EQ(CL_SUCCESS, retVal);
@ -1160,6 +1189,9 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenKernelWithExecEnvIsUsedThen
retVal = clReleaseProgram(pProgram);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = clReleaseCommandQueue(cmdQ);
EXPECT_EQ(CL_SUCCESS, retVal);
retVal = clReleaseContext(context);
EXPECT_EQ(CL_SUCCESS, retVal);
}
@ -2118,6 +2150,24 @@ TEST_F(GTPinTests, givenInitializedGTPinInterfaceWhenLowMemoryConditionOccursThe
injectFailures(allocBufferFunc);
}
TEST_F(GTPinTests, givenParentKernelWhenGtPinAddingSurfaceStateThenItIsNotAddedAndFalseIsReturned) {
GFXCORE_FAMILY genFamily = pDevice->getHardwareInfo().platform.eRenderCoreFamily;
GTPinHwHelper &gtpinHelper = GTPinHwHelper::get(genFamily);
std::unique_ptr<MockParentKernel> parentKernel(MockParentKernel::create(*pContext));
parentKernel->mockKernelInfo->usesSsh = true;
parentKernel->sshLocalSize = 64;
parentKernel->pSshLocal.reset(new char[64]);
size_t sizeSurfaceStates1 = parentKernel->getSurfaceStateHeapSize();
bool surfaceAdded = gtpinHelper.addSurfaceState(parentKernel.get());
EXPECT_FALSE(surfaceAdded);
size_t sizeSurfaceStates2 = parentKernel->getSurfaceStateHeapSize();
EXPECT_EQ(sizeSurfaceStates2, sizeSurfaceStates1);
}
TEST_F(GTPinTests, givenKernelWithSSHThenVerifyThatSSHResizeWorksWell) {
cl_kernel kernel = nullptr;
cl_program pProgram = nullptr;

4
opencl/test/unit_test/mocks/mock_kernel.h
View File

@ -356,7 +356,11 @@ class MockKernelWithInternals {
class MockParentKernel : public Kernel {
public:
using Kernel::auxTranslationRequired;
using Kernel::kernelInfo;
using Kernel::patchBlocksCurbeWithConstantValues;
using Kernel::pSshLocal;
using Kernel::sshLocalSize;
static MockParentKernel *create(Context &context, bool addChildSimdSize = false, bool addChildGlobalMemory = false, bool addChildConstantMemory = false, bool addPrintfForParent = true, bool addPrintfForBlock = true) {
Device &device = context.getDevice(0)->getDevice();