[SYCL] Improve error handling for kernel invocation

sycl/include/CL/sycl/detail/pi.h

@@ -61,7 +61,7 @@ using pi_bitfield = pi_uint64;
 //
 typedef enum {
   PI_SUCCESS = CL_SUCCESS,
-  PI_RESULT_INVALID_KERNEL_NAME = CL_INVALID_KERNEL_NAME,
+  PI_INVALID_KERNEL_NAME = CL_INVALID_KERNEL_NAME,
   PI_INVALID_OPERATION = CL_INVALID_OPERATION,
   PI_INVALID_KERNEL = CL_INVALID_KERNEL,
   PI_INVALID_QUEUE_PROPERTIES = CL_INVALID_QUEUE_PROPERTIES,
@@ -83,6 +83,11 @@ typedef enum {
   PI_COMPILER_NOT_AVAILABLE = CL_COMPILER_NOT_AVAILABLE,
   PI_PROFILING_INFO_NOT_AVAILABLE = CL_PROFILING_INFO_NOT_AVAILABLE,
   PI_DEVICE_NOT_FOUND = CL_DEVICE_NOT_FOUND,
+  PI_INVALID_WORK_ITEM_SIZE = CL_INVALID_WORK_ITEM_SIZE,
+  PI_INVALID_KERNEL_ARGS = CL_INVALID_KERNEL_ARGS,
+  PI_INVALID_IMAGE_SIZE = CL_INVALID_IMAGE_SIZE,
+  PI_IMAGE_FORMAT_NOT_SUPPORTED = CL_IMAGE_FORMAT_NOT_SUPPORTED,
+  PI_MEM_OBJECT_ALLOCATION_FAILURE = CL_MEM_OBJECT_ALLOCATION_FAILURE,
   PI_ERROR_UNKNOWN = -999
 } _pi_result;
 
@@ -98,7 +103,7 @@ typedef enum {
   PI_PLATFORM_INFO_NAME = CL_PLATFORM_NAME,
   PI_PLATFORM_INFO_PROFILE = CL_PLATFORM_PROFILE,
   PI_PLATFORM_INFO_VENDOR = CL_PLATFORM_VENDOR,
-  PI_PLATFORM_INFO_VERSION = CL_PLATFORM_VERSION,
+  PI_PLATFORM_INFO_VERSION = CL_PLATFORM_VERSION
 } _pi_platform_info;
 
 typedef enum {

sycl/source/detail/error_handling/enqueue_kernel.cpp

@@ -166,6 +166,7 @@ bool oclHandleInvalidWorkGroupSize(const device_impl &DeviceImpl,
   // program source.
 
   // Fallback
+
   constexpr pi_result Error = PI_INVALID_WORK_GROUP_SIZE;
   throw runtime_error(
       "OpenCL API failed. OpenCL API returns: " + codeToString(Error), Error);
@@ -177,7 +178,6 @@ bool handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
 
   const plugin &Plugin = DeviceImpl.getPlugin();
   RT::PiDevice Device = DeviceImpl.getHandleRef();
-
   if (HasLocalSize) {
     size_t MaxThreadsPerBlock[3] = {};
     Plugin.call<PiApiKind::piDeviceGetInfo>(
@@ -211,14 +211,76 @@ bool handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
       "PI backend failed. PI backend returns: " + codeToString(Error), Error);
 }
 
+bool handleInvalidWorkItemSize(const device_impl &DeviceImpl,
+                               const NDRDescT &NDRDesc) {
+
+  const plugin &Plugin = DeviceImpl.getPlugin();
+  RT::PiDevice Device = DeviceImpl.getHandleRef();
+
+  size_t MaxWISize[] = {0, 0, 0};
+
+  Plugin.call<PiApiKind::piDeviceGetInfo>(
+      Device, PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES, sizeof(MaxWISize), &MaxWISize,
+      nullptr);
+  for (int i = 0; i < NDRDesc.Dims; i++) {
+    if (NDRDesc.LocalSize[i] > MaxWISize[i])
+      throw sycl::nd_range_error(
+          "Number of work-items in a work-group exceed limit for dimension "
+          "{I}: {NDRDesc.LocalSize[I]} > {MaxWISize[I]}",
+          PI_INVALID_WORK_ITEM_SIZE);
+  }
+  return 0;
+}
+
 bool handleError(pi_result Error, const device_impl &DeviceImpl,
                  pi_kernel Kernel, const NDRDescT &NDRDesc) {
   assert(Error != PI_SUCCESS &&
          "Success is expected to be handled on caller side");
   switch (Error) {
   case PI_INVALID_WORK_GROUP_SIZE:
     return handleInvalidWorkGroupSize(DeviceImpl, Kernel, NDRDesc);
-  // TODO: Handle other error codes
+
+  case PI_INVALID_KERNEL_ARGS:
+    throw sycl::nd_range_error(
+        "The kernel argument values have not been specified "
+        " OR "
+        "a kernel argument declared to be a pointer to a type.",
+        PI_INVALID_KERNEL_ARGS);
+
+  case PI_INVALID_WORK_ITEM_SIZE:
+    return handleInvalidWorkItemSize(DeviceImpl, NDRDesc);
+
+  case PI_IMAGE_FORMAT_NOT_SUPPORTED:
+    throw sycl::nd_range_error(
+        "image object is specified as an argument value"
+        " and the image format is not supported by device associated"
+        " with queue",
+        PI_IMAGE_FORMAT_NOT_SUPPORTED);
+
+  case PI_MISALIGNED_SUB_BUFFER_OFFSET:
+    throw sycl::nd_range_error(
+        "a sub-buffer object is specified as the value for an argument "
+        " that is a buffer object and the offset specified "
+        "when the sub-buffer object is created is not aligned "
+        "to CL_DEVICE_MEM_BASE_ADDR_ALIGN value for device associated"
+        " with queue",
+        PI_MISALIGNED_SUB_BUFFER_OFFSET);
+
+  case PI_MEM_OBJECT_ALLOCATION_FAILURE:
+    throw sycl::nd_range_error(
+        "failure to allocate memory for data store associated with image"
+        " or buffer objects specified as arguments to kernel",
+        PI_MEM_OBJECT_ALLOCATION_FAILURE);
+
+  case PI_INVALID_IMAGE_SIZE:
+    throw sycl::nd_range_error(
+        "image object is specified as an argument value and the image "
+        "dimensions (image width, height, specified or compute row and/or "
+        "slice pitch) are not supported by device associated with queue",
+        PI_INVALID_IMAGE_SIZE);
+
+    // TODO: Handle other error codes
+
   default:
     throw runtime_error(
         "OpenCL API failed. OpenCL API returns: " + codeToString(Error), Error);

sycl/source/detail/program_impl.cpp

@@ -380,7 +380,7 @@ RT::PiKernel program_impl::get_pi_kernel(const string_class &KernelName) const {
     const detail::plugin &Plugin = getPlugin();
     RT::PiResult Err = Plugin.call_nocheck<PiApiKind::piKernelCreate>(
         MProgram, KernelName.c_str(), &Kernel);
-    if (Err == PI_RESULT_INVALID_KERNEL_NAME) {
+    if (Err == PI_INVALID_KERNEL_NAME) {
       throw invalid_object_error(
           "This instance of program does not contain the kernel requested",
           Err);

sycl/source/detail/program_manager/program_manager.cpp

@@ -909,7 +909,7 @@ ProgramManager::getKernelSetId(OSModuleHandle M,
     return ModuleKSIdIt->second;
 
   throw runtime_error("No kernel named " + KernelName + " was found",
-                      PI_RESULT_INVALID_KERNEL_NAME);
+                      PI_INVALID_KERNEL_NAME);
 }
 
 RT::PiDeviceBinaryType ProgramManager::getFormat(const DeviceImage &Img) const {