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

Error from clEnqueueNDRangeKernel() for too big group counts

Browse Source 
Resolves: NEO-6976

Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>
This commit is contained in:
Mateusz Hoppe
2022-10-04 16:57:51 +00:00
committed by Compute-Runtime-Automation
parent 322719e7a2
commit 2f24ef6855
2 changed files with 39 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
opencl/source/command_queue/enqueue_kernel.h
View File

@@ -131,6 +131,13 @@ cl_int CommandQueueHw<GfxFamily>::enqueueKernel(
return CL_INVALID_WORK_GROUP_SIZE; return CL_INVALID_WORK_GROUP_SIZE;
} }
for (auto i = 0u; i < workDim; i++) {
uint64_t dimension = static_cast<uint64_t>(region[i]) / workGroupSize[i];
if (dimension > std::numeric_limits<uint32_t>::max()) {
return CL_INVALID_GLOBAL_WORK_SIZE;
}
}
return enqueueHandler<CL_COMMAND_NDRANGE_KERNEL>( return enqueueHandler<CL_COMMAND_NDRANGE_KERNEL>(
surfaces, surfaces,
false, false,

32
opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
View File

@@ -1502,6 +1502,38 @@ TEST_F(EnqueueKernelTest, givenEnqueueCommandWithWorkDimLargerThanAllowedWhenEnq
EXPECT_EQ(CL_INVALID_WORK_DIMENSION, status); EXPECT_EQ(CL_INVALID_WORK_DIMENSION, status);
} }
TEST_F(EnqueueKernelTest, givenEnqueueCommandWithWorkDimsResultingInMoreThan32BitMaxGroupsWhenEnqueueNDRangeKernelIsCalledThenInvalidGlobalSizeIsReturned) {
if (sizeof(size_t) < 8) {
GTEST_SKIP();
}
size_t max32Bit = std::numeric_limits<uint32_t>::max();
size_t globalWorkSize[3] = {max32Bit * 4, 4, 4};
size_t localWorkSize[3] = {4, 4, 4};
MockKernelWithInternals mockKernel(*pClDevice);
auto testedWorkDim = 3;
auto status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, status);
globalWorkSize[0] = max32Bit * 4 + 4;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
globalWorkSize[0] = 4;
globalWorkSize[1] = max32Bit * 4 + 4;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
globalWorkSize[1] = 4;
globalWorkSize[2] = max32Bit * 4 + 4;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, localWorkSize, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
}
HWTEST_F(EnqueueKernelTest, givenVMEKernelWhenEnqueueKernelThenDispatchFlagsHaveMediaSamplerRequired) { HWTEST_F(EnqueueKernelTest, givenVMEKernelWhenEnqueueKernelThenDispatchFlagsHaveMediaSamplerRequired) {
auto mockCsr = new MockCsrHw2<FamilyType>(*pDevice->executionEnvironment, pDevice->getRootDeviceIndex(), pDevice->getDeviceBitfield()); auto mockCsr = new MockCsrHw2<FamilyType>(*pDevice->executionEnvironment, pDevice->getRootDeviceIndex(), pDevice->getDeviceBitfield());
mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch); mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch);