diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h
index 507f15fbd03ad..75cde167f0c65 100644
--- a/sycl/include/CL/sycl/detail/pi.h
+++ b/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 {
diff --git a/sycl/source/detail/error_handling/enqueue_kernel.cpp b/sycl/source/detail/error_handling/enqueue_kernel.cpp
index ded4d1e6c8e71..f73c77b1d7649 100644
--- a/sycl/source/detail/error_handling/enqueue_kernel.cpp
+++ b/sycl/source/detail/error_handling/enqueue_kernel.cpp
@@ -211,6 +211,28 @@ 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 (unsigned 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 " +
+              std::to_string(I) + " : " + std::to_string(NDRDesc.LocalSize[I]) +
+              " > " + std::to_string(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 &&
@@ -218,7 +240,48 @@ bool handleError(pi_result Error, const device_impl &DeviceImpl,
   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);
diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp
index a59f2414e11d0..9785e689fd600 100644
--- a/sycl/source/detail/program_impl.cpp
+++ b/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);
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index d9719547c7a0c..9e676045e8d08 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/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 {