OCL: Fix error for too big dimensions

- 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>

Source: f77f470

Author: HoppeMateusz

opencl/source/command_queue/enqueue_common.h

@@ -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);
     }

opencl/source/command_queue/enqueue_kernel.h

@@ -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,

opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp

@@ -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);