[SYCL] Improve diagnostics for invalid work group size with L0 backend. (#8378)

Prior to this change, when requesting more work items than available
while using L0 as backend, the error message was quite vague and didn't
really help users understand the problem.

---------

Signed-off-by: Maronas, Marcos <[email protected]>

Author: maarquitos14

sycl/source/detail/error_handling/enqueue_kernel.cpp

@@ -51,12 +51,17 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
   // versions.  If this is an OpenCL backend, get the version.
   bool IsOpenCL = false;    // Backend is any OpenCL version
   bool IsOpenCLV1x = false; // Backend is OpenCL 1.x
-  bool IsOpenCLV20 = false; // Backend is OpenCL 2.0
-  if (Platform.get_backend() == sycl::backend::opencl) {
+  bool IsOpenCLVGE20 = false; // Backend is Greater or Equal to OpenCL 2.0
+  bool IsLevelZero = false;   // Backend is any OneAPI Level 0 version
+  auto Backend = Platform.get_backend();
+  if (Backend == sycl::backend::opencl) {
     std::string VersionString = DeviceImpl.get_info<info::device::version>();
     IsOpenCL = true;
     IsOpenCLV1x = (VersionString.find("1.") == 0);
-    IsOpenCLV20 = (VersionString.find("2.0") == 0);
+    IsOpenCLVGE20 =
+        (VersionString.find("2.") == 0) || (VersionString.find("3.") == 0);
+  } else if (Backend == sycl::backend::ext_oneapi_level_zero) {
+    IsLevelZero = true;
   }
 
   size_t CompileWGSize[3] = {0};
@@ -69,7 +74,7 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
     // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is NULL and the
     // reqd_work_group_size attribute is used to declare the work-group size
     // for kernel in the program source.
-    if (!HasLocalSize && (IsOpenCLV1x || IsOpenCLV20)) {
+    if (!HasLocalSize && (IsOpenCLV1x || IsOpenCLVGE20)) {
       throw sycl::nd_range_error(
           "OpenCL 1.x and 2.0 requires to pass local size argument even if "
           "required work-group size was specified in the program source",
@@ -91,7 +96,6 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
               std::to_string(CompileWGSize[0]) + "}",
           PI_ERROR_INVALID_WORK_GROUP_SIZE);
   }
-  if (IsOpenCL) {
     if (IsOpenCLV1x) {
       // OpenCL 1.x:
       // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and
@@ -110,8 +114,8 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
             "Total number of work-items in a work-group cannot exceed " +
                 std::to_string(MaxWGSize),
             PI_ERROR_INVALID_WORK_GROUP_SIZE);
-    } else {
-      // OpenCL 2.x:
+    } else if (IsOpenCLVGE20 || IsLevelZero) {
+      // OpenCL 2.x or OneAPI Level Zero:
       // PI_ERROR_INVALID_WORK_GROUP_SIZE if local_work_size is specified and
       // the total number of work-items in the work-group computed as
       // local_work_size[0] * ... * local_work_size[work_dim - 1] is greater
@@ -128,10 +132,9 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
             "Total number of work-items in a work-group cannot exceed " +
                 std::to_string(KernelWGSize) + " for this kernel",
             PI_ERROR_INVALID_WORK_GROUP_SIZE);
+    } else {
+      // TODO: Should probably have something similar for the other backends
     }
-  } else {
-    // TODO: Should probably have something similar for the other backends
-  }
 
   if (HasLocalSize) {
     // Is the global range size evenly divisible by the local workgroup size?