Media/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2025-09-10 12:53:42 +08:00
Code Issues Packages Projects Releases Wiki Activity

Add error handling when kernel uses too much SLM

Browse Source 
API Functions reporting error:
- clCreateKernel
- clEnqueueNDRangeKernel
- zeKernelCreate
- zeCommandListAppendLaunchKernel
- zeKernelSuggestGroupSize

Related-To: NEO-7280
Signed-off-by: Fabian Zwolinski <fabian.zwolinski@intel.com>
This commit is contained in:
Fabian Zwolinski
2022-09-08 12:01:51 +00:00
committed by Compute-Runtime-Automation
parent 5a0dd8048b
commit 1142404c0c
11 changed files with 199 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

42
opencl/test/unit_test/api/cl_enqueue_nd_range_kernel_tests.inl
View File

@ -1,5 +1,5 @@
/*
* Copyright (C) 2018-2021 Intel Corporation
* Copyright (C) 2018-2022 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
@ -39,6 +39,46 @@ TEST_F(clEnqueueNDRangeKernelTests, GivenValidParametersWhenExecutingKernelThenS
EXPECT_EQ(CL_SUCCESS, retVal);
}
TEST_F(clEnqueueNDRangeKernelTests, GivenKernelWithSlmSizeExceedingLocalMemorySizeWhenExecutingKernelThenOutOfResourcesIsReturned) {
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};
cl_uint numEventsInWaitList = 0;
cl_event *eventWaitList = nullptr;
cl_event *event = nullptr;
auto localMemSize = static_cast<uint32_t>(pDevice->getDevice().getDeviceInfo().localMemSize);
pProgram->mockKernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize - 10u;
retVal = clEnqueueNDRangeKernel(
pCommandQueue,
pMultiDeviceKernel,
workDim,
globalWorkOffset,
globalWorkSize,
localWorkSize,
numEventsInWaitList,
eventWaitList,
event);
EXPECT_EQ(CL_SUCCESS, retVal);
pProgram->mockKernelInfo.kernelDescriptor.kernelAttributes.slmInlineSize = localMemSize + 10u;
retVal = clEnqueueNDRangeKernel(
pCommandQueue,
pMultiDeviceKernel,
workDim,
globalWorkOffset,
globalWorkSize,
localWorkSize,
numEventsInWaitList,
eventWaitList,
event);
EXPECT_EQ(CL_OUT_OF_RESOURCES, retVal);
}
TEST_F(clEnqueueNDRangeKernelTests, GivenQueueIncapableWhenExecutingKernelThenInvalidOperationIsReturned) {
cl_uint workDim = 1;
size_t globalWorkOffset[3] = {0, 0, 0};