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

OCL: Fix error for too big dimensions

Browse Source 
- cases with null lws should only fail when computed
lws sizes result in too big number of workgroups

Related-To: NEO-6976

Signed-off-by: Mateusz Hoppe <mateusz.hoppe@intel.com>
This commit is contained in:
Mateusz Hoppe
2022-11-08 13:33:22 +00:00
committed by Compute-Runtime-Automation
parent 7be9881758
commit f77f47099c
3 changed files with 74 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
opencl/source/command_queue/enqueue_common.h
View File

@ -94,6 +94,21 @@ cl_int CommandQueueHw<GfxFamily>::enqueueHandler(Surface *(&surfaces)[surfaceCou
}
}
if (commandType == CL_COMMAND_NDRANGE_KERNEL) {
if (!multiDispatchInfo.empty()) {
for (auto &dispatchInfo : multiDispatchInfo) {
auto nwgs = dispatchInfo.getNumberOfWorkgroups();
for (auto i = 0u; i < workDim; i++) {
uint64_t dimension = static_cast<uint64_t>(nwgs[i]);
if (dimension > std::numeric_limits<uint32_t>::max()) {
return CL_INVALID_GLOBAL_WORK_SIZE;
}
}
}
}
}
if (AuxTranslationMode::Builtin == auxTranslationMode) {
dispatchAuxTranslationBuiltin(multiDispatchInfo, AuxTranslationDirection::NonAuxToAux);
}

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

@ -131,13 +131,6 @@ cl_int CommandQueueHw<GfxFamily>::enqueueKernel(
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>(
surfaces,
false,

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

@ -1534,6 +1534,65 @@ TEST_F(EnqueueKernelTest, givenEnqueueCommandWithWorkDimsResultingInMoreThan32Bi
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
}
TEST_F(EnqueueKernelTest, givenEnqueueCommandWithNullLwsAndWorkDimsResultingInMoreThan32BitMaxGroupsWhenEnqueueNDRangeKernelIsCalledThenInvalidGlobalSizeIsReturned) {
if (sizeof(size_t) < 8) {
GTEST_SKIP();
}
auto maxWgSize = static_cast<uint32_t>(pClDevice->getDevice().getDeviceInfo().maxWorkGroupSize);
size_t max32Bit = std::numeric_limits<uint32_t>::max();
size_t globalWorkSize[3] = {(max32Bit + 1) * maxWgSize, 3, 4};
MockKernelWithInternals mockKernel(*pClDevice);
auto testedWorkDim = 3;
auto status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
globalWorkSize[0] = (max32Bit + 1) * maxWgSize + 3;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
globalWorkSize[0] = 4;
globalWorkSize[1] = (max32Bit + 1) * maxWgSize;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
globalWorkSize[1] = 4;
globalWorkSize[2] = (max32Bit + 1) * maxWgSize * 2 + 3;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_INVALID_GLOBAL_WORK_SIZE, status);
}
TEST_F(EnqueueKernelTest, givenEnqueueCommandWithNullLwsAndWorkDimsResultingInLessThan32BitMaxGroupsWhenEnqueueNDRangeKernelIsCalledThenSuccessIsReturned) {
if (sizeof(size_t) < 8) {
GTEST_SKIP();
}
size_t max32Bit = std::numeric_limits<uint32_t>::max();
size_t globalWorkSize[3] = {(max32Bit + 1) * 4, 1, 1};
MockKernelWithInternals mockKernel(*pClDevice);
auto testedWorkDim = 3;
auto status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, status);
globalWorkSize[0] = 1;
globalWorkSize[1] = (max32Bit + 1) * 4;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, status);
globalWorkSize[1] = 1;
globalWorkSize[2] = (max32Bit + 1) * 4;
status = clEnqueueNDRangeKernel(pCmdQ, mockKernel.mockMultiDeviceKernel, testedWorkDim, nullptr, globalWorkSize, nullptr, 0, nullptr, nullptr);
EXPECT_EQ(CL_SUCCESS, status);
}
HWTEST_F(EnqueueKernelTest, givenVMEKernelWhenEnqueueKernelThenDispatchFlagsHaveMediaSamplerRequired) {
auto mockCsr = new MockCsrHw2<FamilyType>(*pDevice->executionEnvironment, pDevice->getRootDeviceIndex(), pDevice->getDeviceBitfield());
mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch);