diff --git a/opencl/source/command_queue/enqueue_kernel.h b/opencl/source/command_queue/enqueue_kernel.h
index d911c43605..06edcf786c 100644
--- a/opencl/source/command_queue/enqueue_kernel.h
+++ b/opencl/source/command_queue/enqueue_kernel.h
@@ -131,6 +131,13 @@ 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,
diff --git a/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp b/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
index e5dbeeca38..ffa1958c65 100644
--- a/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
+++ b/opencl/test/unit_test/command_queue/enqueue_kernel_1_tests.cpp
@@ -1502,6 +1502,38 @@ TEST_F(EnqueueKernelTest, givenEnqueueCommandWithWorkDimLargerThanAllowedWhenEnq
     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) {
     auto mockCsr = new MockCsrHw2<FamilyType>(*pDevice->executionEnvironment, pDevice->getRootDeviceIndex(), pDevice->getDeviceBitfield());
     mockCsr->overrideDispatchPolicy(DispatchMode::BatchedDispatch);