diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 74bb361bb96f2..baa869fc3eaf3 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -336,7 +336,8 @@ class image_accessor
   template <info::device param>
   void checkDeviceFeatureSupported(const device &Device) {
     if (!Device.get_info<param>())
-      throw feature_not_supported("Images are not supported by this device.");
+      throw feature_not_supported("Images are not supported by this device.",
+                                  PI_INVALID_OPERATION);
   }
 
 #ifdef __SYCL_DEVICE_ONLY__
@@ -357,8 +358,8 @@ class image_accessor
 
   sycl::vec<int, Dimensions> getRangeInternal() const {
     // TODO: Implement for host.
-    throw runtime_error(
-        "image::getRangeInternal() is not implemented for host");
+    throw runtime_error("image::getRangeInternal() is not implemented for host",
+                        PI_INVALID_OPERATION);
     return sycl::vec<int, Dimensions>{1};
   }
 
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index b894bfeda9961..cc19eef9257fe 100644
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -180,13 +180,14 @@ class buffer {
         IsSubBuffer(true) {
     if (b.is_sub_buffer())
       throw cl::sycl::invalid_object_error(
-          "Cannot create sub buffer from sub buffer.");
+          "Cannot create sub buffer from sub buffer.", PI_INVALID_VALUE);
     if (isOutOfBounds(baseIndex, subRange, b.Range))
       throw cl::sycl::invalid_object_error(
-          "Requested sub-buffer size exceeds the size of the parent buffer");
+          "Requested sub-buffer size exceeds the size of the parent buffer",
+          PI_INVALID_VALUE);
     if (!isContiguousRegion(baseIndex, subRange, b.Range))
       throw cl::sycl::invalid_object_error(
-          "Requested sub-buffer region is not contiguous");
+          "Requested sub-buffer region is not contiguous", PI_INVALID_VALUE);
   }
 
   template <int N = dimensions, typename = EnableIfOneDimension<N>>
@@ -282,7 +283,8 @@ class buffer {
       throw cl::sycl::invalid_object_error(
           "Total size in bytes represented by the type and range of the "
           "reinterpreted SYCL buffer does not equal the total size in bytes "
-          "represented by the type and range of this SYCL buffer");
+          "represented by the type and range of this SYCL buffer",
+          PI_INVALID_VALUE);
 
     return buffer<ReinterpretT, ReinterpretDim, AllocatorT>(
         impl, reinterpretRange, OffsetInBytes, IsSubBuffer);
diff --git a/sycl/include/CL/sycl/detail/array.hpp b/sycl/include/CL/sycl/detail/array.hpp
index 1a4eaa16820d4..d4aada6d7dd86 100644
--- a/sycl/include/CL/sycl/detail/array.hpp
+++ b/sycl/include/CL/sycl/detail/array.hpp
@@ -111,7 +111,8 @@ template <int dimensions = 1> class array {
   ALWAYS_INLINE void check_dimension(int dimension) const {
 #ifndef __SYCL_DEVICE_ONLY__
     if (dimension >= dimensions || dimension < 0) {
-      throw cl::sycl::invalid_parameter_error("Index out of range");
+      throw cl::sycl::invalid_parameter_error("Index out of range",
+                                              PI_INVALID_VALUE);
     }
 #endif
   }
diff --git a/sycl/include/CL/sycl/detail/cg.hpp b/sycl/include/CL/sycl/detail/cg.hpp
index c45e00643576c..e4d0a0d30967c 100644
--- a/sycl/include/CL/sycl/detail/cg.hpp
+++ b/sycl/include/CL/sycl/detail/cg.hpp
@@ -279,7 +279,8 @@ class HostKernel : public HostKernelBase {
     for (int I = 0; I < Dims; ++I) {
       if (NDRDesc.LocalSize[I] == 0 ||
           NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
-        throw sycl::nd_range_error("Invalid local size for global size");
+        throw sycl::nd_range_error("Invalid local size for global size",
+                                   PI_INVALID_WORK_GROUP_SIZE);
       GroupSize[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
     }
 
@@ -320,7 +321,8 @@ class HostKernel : public HostKernelBase {
     for (int I = 0; I < Dims; ++I) {
       if (NDRDesc.LocalSize[I] == 0 ||
           NDRDesc.GlobalSize[I] % NDRDesc.LocalSize[I] != 0)
-        throw sycl::nd_range_error("Invalid local size for global size");
+        throw sycl::nd_range_error("Invalid local size for global size",
+                                   PI_INVALID_WORK_GROUP_SIZE);
       NGroups[I] = NDRDesc.GlobalSize[I] / NDRDesc.LocalSize[I];
     }
 
diff --git a/sycl/include/CL/sycl/detail/image_accessor_util.hpp b/sycl/include/CL/sycl/detail/image_accessor_util.hpp
index 9d3375f98a453..f0a4cefbe43b0 100644
--- a/sycl/include/CL/sycl/detail/image_accessor_util.hpp
+++ b/sycl/include/CL/sycl/detail/image_accessor_util.hpp
@@ -193,7 +193,8 @@ vec<T, 4> readPixel(T *Ptr, const image_channel_order ChannelOrder,
     Pixel.x() = Ptr[3]; // r
     break;
   default:
-    throw cl::sycl::invalid_parameter_error("Unhandled image channel order");
+    throw cl::sycl::invalid_parameter_error("Unhandled image channel order",
+                                            PI_INVALID_VALUE);
   }
 
   return Pixel;
@@ -265,7 +266,8 @@ void writePixel(const vec<T, 4> Pixel, T *Ptr,
     Ptr[3] = Pixel.x(); // r
     break;
   default:
-    throw cl::sycl::invalid_parameter_error("Unhandled image channel order");
+    throw cl::sycl::invalid_parameter_error("Unhandled image channel order",
+                                            PI_INVALID_VALUE);
   }
 }
 
@@ -293,7 +295,8 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
     // unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of read data - cl_uint4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   }
 }
 
@@ -314,7 +317,8 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
     // signed_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of read data - cl_int4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   }
 }
 
@@ -395,7 +399,8 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
     // and signed/unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of read data - cl_float4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   case image_channel_type::fp16:
     // Host has conversion from float to half with accuracy as required in
     // section 8.3.2 OpenCL spec.
@@ -425,7 +430,8 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
     // TODO: Missing information in OpenCL spec.
     throw cl::sycl::feature_not_supported(
         "Currently unsupported datatype conversion from image_channel_type "
-        "to cl_half4.");
+        "to cl_half4.",
+        PI_INVALID_OPERATION);
   case image_channel_type::signed_int8:
   case image_channel_type::signed_int16:
   case image_channel_type::signed_int32:
@@ -437,14 +443,16 @@ void convertReadData(const vec<ChannelType, 4> PixelData,
     // and signed/unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype to read- cl_half4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   case image_channel_type::fp16:
     RetData = PixelData.template convert<cl_half>();
     break;
   case image_channel_type::fp32:
     throw cl::sycl::invalid_parameter_error(
         "Datatype to read - cl_half4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   default:
     break;
   }
@@ -484,7 +492,8 @@ convertWriteData(const vec<cl_uint, 4> WriteData,
     // unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of data to write - cl_uint4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   }
 }
 
@@ -516,7 +525,8 @@ convertWriteData(const vec<cl_int, 4> WriteData,
     // signed_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of data to write - cl_int4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   }
 }
 
@@ -554,7 +564,8 @@ convertWriteData(const vec<cl_float, 4> WriteData,
     // TODO: Missing information in OpenCL spec.
     throw cl::sycl::feature_not_supported(
         "Currently unsupported datatype conversion from image_channel_type "
-        "to cl_float4.");
+        "to cl_float4.",
+        PI_INVALID_OPERATION);
   case image_channel_type::unorm_short_555:
     // TODO: Missing information in OpenCL spec.
     // Check if the below code is correct after the spec is updated.
@@ -596,7 +607,8 @@ convertWriteData(const vec<cl_float, 4> WriteData,
     // and signed/unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of data to write - cl_float4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   case image_channel_type::fp16:
     // Host has conversion from float to half with accuracy as required in
     // section 8.3.2 OpenCL spec.
@@ -624,7 +636,8 @@ convertWriteData(const vec<cl_half, 4> WriteData,
     // TODO: Missing information in OpenCL spec.
     throw cl::sycl::feature_not_supported(
         "Currently unsupported datatype conversion from image_channel_type "
-        "to cl_half4.");
+        "to cl_half4.",
+        PI_INVALID_OPERATION);
   case image_channel_type::signed_int8:
   case image_channel_type::signed_int16:
   case image_channel_type::signed_int32:
@@ -636,13 +649,15 @@ convertWriteData(const vec<cl_half, 4> WriteData,
     // and signed/unsigned_int32.
     throw cl::sycl::invalid_parameter_error(
         "Datatype of data to write - cl_float4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   case image_channel_type::fp16:
     return WriteData.convert<ChannelType>();
   case image_channel_type::fp32:
     throw cl::sycl::invalid_parameter_error(
         "Datatype of data to write - cl_float4 is incompatible with the "
-        "image_channel_type of the image.");
+        "image_channel_type of the image.",
+        PI_INVALID_VALUE);
   default:
     break;
   }
@@ -1007,7 +1022,8 @@ DataT imageReadSamplerHostImpl(const CoordT &Coords, const sampler &Smpl,
       throw cl::sycl::feature_not_supported(
           "Sampler used with unsupported configuration of "
           "mirrored_repeat/repeat filtering mode with unnormalized "
-          "coordinates. ");
+          "coordinates. ",
+          PI_INVALID_OPERATION);
     case addressing_mode::clamp_to_edge:
     case addressing_mode::clamp:
     case addressing_mode::none:
diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h
index ea0f5d5703eca..ff72923551723 100644
--- a/sycl/include/CL/sycl/detail/pi.h
+++ b/sycl/include/CL/sycl/detail/pi.h
@@ -73,6 +73,9 @@ typedef enum {
   PI_MISALIGNED_SUB_BUFFER_OFFSET = CL_MISALIGNED_SUB_BUFFER_OFFSET,
   PI_BUILD_PROGRAM_FAILURE = CL_BUILD_PROGRAM_FAILURE,
   PI_INVALID_WORK_GROUP_SIZE = CL_INVALID_WORK_GROUP_SIZE,
+  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_ERROR_UNKNOWN = -999
 } _pi_result;
 
diff --git a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
index 940210a7a76cd..4bd7641953099 100644
--- a/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
+++ b/sycl/include/CL/sycl/detail/sycl_mem_obj_t.hpp
@@ -234,9 +234,10 @@ class SYCLMemObjT : public SYCLMemObjI {
     MHostPtrReadOnly = iterator_to_const_type_t<InputIterator>::value;
     setAlign(RequiredAlign);
     if (useHostPtr())
-      throw invalid_parameter_error(
+      throw runtime_error(
           "Buffer constructor from a pair of iterator values does not support "
-          "use_host_ptr property.");
+          "use_host_ptr property.",
+          PI_INVALID_OPERATION);
 
     setAlign(RequiredAlign);
     MShadowCopy = allocateHostMem();
diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp
index 826f394e99b0d..781d108ec0dfa 100644
--- a/sycl/include/CL/sycl/exception.hpp
+++ b/sycl/include/CL/sycl/exception.hpp
@@ -11,6 +11,7 @@
 // 4.9.2 Exception Class Interface
 
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/detail/pi.h>
 #include <CL/sycl/stl.hpp>
 
 #include <exception>
@@ -37,15 +38,15 @@ class exception: public std::exception {
 
 private:
   string_class MMsg;
-  cl_int MCLErr = CL_SUCCESS;
+  cl_int MCLErr;
   shared_ptr_class<context> MContext;
 
 protected:
-  exception(const char *Msg, const cl_int CLErr = CL_SUCCESS,
+  exception(const char *Msg, const cl_int CLErr,
             shared_ptr_class<context> Context = nullptr)
       : exception(string_class(Msg), CLErr, Context) {}
 
-  exception(const string_class &Msg, const cl_int CLErr = CL_SUCCESS,
+  exception(const string_class &Msg, const cl_int CLErr,
             shared_ptr_class<context> Context = nullptr)
       : MMsg(Msg + " " + detail::codeToString(CLErr)), MCLErr(CLErr),
         MContext(Context) {}
@@ -55,11 +56,10 @@ class runtime_error : public exception {
 public:
   runtime_error() = default;
 
-  runtime_error(const char *Msg, cl_int Err = CL_SUCCESS)
+  runtime_error(const char *Msg, cl_int Err)
       : runtime_error(string_class(Msg), Err) {}
 
-  runtime_error(const string_class &Msg, cl_int Err = CL_SUCCESS)
-      : exception(Msg, Err) {}
+  runtime_error(const string_class &Msg, cl_int Err) : exception(Msg, Err) {}
 };
 class kernel_error : public runtime_error {
   using runtime_error::runtime_error;
@@ -80,11 +80,10 @@ class device_error : public exception {
 public:
   device_error() = default;
 
-  device_error(const char *Msg, cl_int Err = CL_SUCCESS)
+  device_error(const char *Msg, cl_int Err)
       : device_error(string_class(Msg), Err) {}
 
-  device_error(const string_class &Msg, cl_int Err = CL_SUCCESS)
-      : exception(Msg, Err) {}
+  device_error(const string_class &Msg, cl_int Err) : exception(Msg, Err) {}
 };
 class compile_program_error : public device_error {
   using device_error::device_error;
diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp
index dd464de3d851c..a753190079f49 100644
--- a/sycl/include/CL/sycl/handler.hpp
+++ b/sycl/include/CL/sycl/handler.hpp
@@ -296,10 +296,12 @@ class handler {
   void verifySyclKernelInvoc(const kernel &SyclKernel) {
     if (is_host()) {
       throw invalid_object_error(
-          "This kernel invocation method cannot be used on the host");
+          "This kernel invocation method cannot be used on the host",
+          PI_INVALID_DEVICE);
     }
     if (SyclKernel.is_host()) {
-      throw invalid_object_error("Invalid kernel type, OpenCL expected");
+      throw invalid_object_error("Invalid kernel type, OpenCL expected",
+                                 PI_INVALID_KERNEL);
     }
   }
 
diff --git a/sycl/include/CL/sycl/intel/function_pointer.hpp b/sycl/include/CL/sycl/intel/function_pointer.hpp
index 3a59bb027aaac..69d32ac0e4940 100644
--- a/sycl/include/CL/sycl/intel/function_pointer.hpp
+++ b/sycl/include/CL/sycl/intel/function_pointer.hpp
@@ -76,7 +76,8 @@ device_func_ptr_holder_t get_device_func_ptr(FuncType F, const char *FuncName,
 
   if (program_state::linked != P.get_state()) {
     throw invalid_parameter_error(
-        "Program must be built before passing to get_device_func_ptr");
+        "Program must be built before passing to get_device_func_ptr",
+        PI_INVALID_OPERATION);
   }
 
   return detail::getDeviceFunctionPointerImpl(D, P, FuncName);
diff --git a/sycl/include/CL/sycl/intel/sub_group_host.hpp b/sycl/include/CL/sycl/intel/sub_group_host.hpp
index 16d9514fa4333..2496d116e28d3 100644
--- a/sycl/include/CL/sycl/intel/sub_group_host.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group_host.hpp
@@ -23,138 +23,166 @@ struct sub_group {
   /* --- common interface members --- */
 
   id<1> get_local_id() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
   range<1> get_local_range() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   range<1> get_max_local_range() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   id<1> get_group_id() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   size_t get_group_range() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   size_t get_uniform_group_range() const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- vote / ballot functions --- */
 
   bool any(bool predicate) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   bool all(bool predicate) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- collectives --- */
 
   template <typename T> T broadcast(T x, id<1> local_id) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T reduce(T x, BinaryOperation op) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T reduce(T x, T init, BinaryOperation op) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T exclusive_scan(T x, BinaryOperation op) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T exclusive_scan(T x, T init, BinaryOperation op) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T inclusive_scan(T x, BinaryOperation op) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, class BinaryOperation>
   T inclusive_scan(T x, BinaryOperation op, T init) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- one - input shuffles --- */
   /* indices in [0 , sub - group size ) */
 
   template <typename T> T shuffle(T x, id<1> local_id) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T> T shuffle_down(T x, uint32_t delta) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
   template <typename T> T shuffle_up(T x, uint32_t delta) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T> T shuffle_xor(T x, id<1> value) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- two - input shuffles --- */
   /* indices in [0 , 2* sub - group size ) */
   template <typename T> T shuffle(T x, T y, id<1> local_id) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
   template <typename T>
   T shuffle_down(T current, T next, uint32_t delta) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
   template <typename T>
   T shuffle_up(T previous, T current, uint32_t delta) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- sub - group load / stores --- */
   /* these can map to SIMD or block read / write hardware where available */
   template <typename T, access::address_space Space>
   T load(const multi_ptr<T, Space> src) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <int N, typename T, access::address_space Space>
   vec<T, N> load(const multi_ptr<T, Space> src) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <typename T, access::address_space Space>
   void store(multi_ptr<T, Space> dst, const T &x) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   template <int N, typename T, access::address_space Space>
   void store(multi_ptr<T, Space> dst, const vec<T, N> &x) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
   /* --- synchronization functions --- */
   void barrier(access::fence_space accessSpace =
                    access::fence_space::global_and_local) const {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 
 protected:
   template <int dimensions> friend class cl::sycl::nd_item;
   sub_group() {
-    throw runtime_error("Subgroups are not supported on host device. ");
+    throw runtime_error("Subgroups are not supported on host device. ",
+                        PI_INVALID_DEVICE);
   }
 };
 } // namespace intel
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
index efc0eb55e1157..472ce566d377a 100644
--- a/sycl/include/CL/sycl/types.hpp
+++ b/sycl/include/CL/sycl/types.hpp
@@ -243,11 +243,12 @@ detail::enable_if_t<is_float_to_int<T, R>::value, R> convertImpl(T Value) {
     int OldRoundingDirection = std::fegetround();
     int Err = std::fesetround(FE_TONEAREST);
     if (Err)
-      throw runtime_error("Unable to set rounding mode to FE_TONEAREST");
+      throw runtime_error("Unable to set rounding mode to FE_TONEAREST",
+                          PI_ERROR_UNKNOWN);
     R Result = std::rint(Value);
     Err = std::fesetround(OldRoundingDirection);
     if (Err)
-      throw runtime_error("Unable to restore rounding mode.");
+      throw runtime_error("Unable to restore rounding mode.", PI_ERROR_UNKNOWN);
     return Result;
   }
     // Round toward zero.
diff --git a/sycl/include/CL/sycl/usm/usm_allocator.hpp b/sycl/include/CL/sycl/usm/usm_allocator.hpp
index ff6526bc1cbf5..d598d4ec49a12 100644
--- a/sycl/include/CL/sycl/usm/usm_allocator.hpp
+++ b/sycl/include/CL/sycl/usm/usm_allocator.hpp
@@ -65,7 +65,8 @@ class usm_allocator {
       typename std::enable_if<AllocT == usm::alloc::device, int>::type = 0>
   void construct(pointer Ptr, const_reference Val) {
     throw feature_not_supported(
-        "Device pointers do not support construct on host");
+        "Device pointers do not support construct on host",
+        PI_INVALID_OPERATION);
   }
 
   /// Destroys an object.
@@ -85,7 +86,7 @@ class usm_allocator {
       typename std::enable_if<AllocT == usm::alloc::device, int>::type = 0>
   void destroy(pointer Ptr) {
     throw feature_not_supported(
-        "Device pointers do not support destroy on host");
+        "Device pointers do not support destroy on host", PI_INVALID_OPERATION);
   }
 
   /// Note:: AllocKind == alloc::device is not allowed.
@@ -104,7 +105,7 @@ class usm_allocator {
       typename std::enable_if<AllocT == usm::alloc::device, int>::type = 0>
   pointer address(reference Val) const {
     throw feature_not_supported(
-        "Device pointers do not support address on host");
+        "Device pointers do not support address on host", PI_INVALID_OPERATION);
   }
 
   template <
@@ -119,7 +120,7 @@ class usm_allocator {
       typename std::enable_if<AllocT == usm::alloc::device, int>::type = 0>
   const_pointer address(const_reference Val) const {
     throw feature_not_supported(
-        "Device pointers do not support address on host");
+        "Device pointers do not support address on host", PI_INVALID_OPERATION);
   }
 
   /// Allocates memory.
diff --git a/sycl/source/context.cpp b/sycl/source/context.cpp
index 1aba28b46a42a..8a4cf83a92520 100644
--- a/sycl/source/context.cpp
+++ b/sycl/source/context.cpp
@@ -40,7 +40,7 @@ context::context(const platform &Platform, async_handler AsyncHandler,
 context::context(const vector_class<device> &DeviceList,
                  async_handler AsyncHandler, bool UsePrimaryContext) {
   if (DeviceList.empty()) {
-    throw invalid_parameter_error("DeviceList is empty.");
+    throw invalid_parameter_error("DeviceList is empty.", PI_INVALID_VALUE);
   }
   auto NonHostDeviceIter = std::find_if_not(
       DeviceList.begin(), DeviceList.end(),
@@ -59,7 +59,8 @@ context::context(const vector_class<device> &DeviceList,
                                  NonHostPlatform));
                     }))
       throw invalid_parameter_error(
-          "Can't add devices across platforms to a single context.");
+          "Can't add devices across platforms to a single context.",
+          PI_INVALID_DEVICE);
     else
       impl = std::make_shared<detail::context_impl>(DeviceList, AsyncHandler,
                                                     UsePrimaryContext);
diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp
index 1eced4ec82ce0..d1bee11d0d8ec 100644
--- a/sycl/source/detail/context_impl.cpp
+++ b/sycl/source/detail/context_impl.cpp
@@ -97,7 +97,8 @@ cl_context context_impl::get() const {
     return pi::cast<cl_context>(MContext);
   }
   throw invalid_object_error(
-      "This instance of context doesn't support OpenCL interoperability.");
+      "This instance of context doesn't support OpenCL interoperability.",
+      PI_INVALID_CONTEXT);
 }
 
 bool context_impl::is_host() const { return MHostContext || !MPluginInterop; }
diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp
index 15c224fe98d48..efce61d181a2d 100644
--- a/sycl/source/detail/device_impl.cpp
+++ b/sycl/source/detail/device_impl.cpp
@@ -72,7 +72,8 @@ bool device_impl::is_affinity_supported(
 
 cl_device_id device_impl::get() const {
   if (MIsHostDevice)
-    throw invalid_object_error("This instance of device is a host instance");
+    throw invalid_object_error("This instance of device is a host instance",
+                               PI_INVALID_DEVICE);
 
   if (!MIsRootDevice) {
     // TODO catch an exception and put it to list of asynchronous exceptions
@@ -137,7 +138,8 @@ device_impl::create_sub_devices(size_t ComputeUnits) const {
   if (MIsHostDevice)
     // TODO: implement host device partitioning
     throw runtime_error(
-        "Partitioning to subdevices of the host device is not implemented yet");
+        "Partitioning to subdevices of the host device is not implemented yet",
+        PI_INVALID_DEVICE);
 
   if (!is_partition_supported(info::partition_property::partition_equally)) {
     throw cl::sycl::feature_not_supported();
@@ -156,7 +158,8 @@ device_impl::create_sub_devices(const vector_class<size_t> &Counts) const {
   if (MIsHostDevice)
     // TODO: implement host device partitioning
     throw runtime_error(
-        "Partitioning to subdevices of the host device is not implemented yet");
+        "Partitioning to subdevices of the host device is not implemented yet",
+        PI_INVALID_DEVICE);
 
   if (!is_partition_supported(
           info::partition_property::partition_by_counts)) {
@@ -176,12 +179,8 @@ vector_class<device> device_impl::create_sub_devices(
   if (MIsHostDevice)
     // TODO: implement host device partitioning
     throw runtime_error(
-        "Partitioning to subdevices of the host device is not implemented yet");
-
-  // TODO: implement host device partitioning
-  if (MIsHostDevice)
-    throw runtime_error(
-        "Partitioning to subdevices of the host device is not implemented yet");
+        "Partitioning to subdevices of the host device is not implemented yet",
+        PI_INVALID_DEVICE);
 
   if (!is_partition_supported(
           info::partition_property::partition_by_affinity_domain) ||
diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp
index ad9aa993a5228..9769d66771978 100644
--- a/sycl/source/detail/device_impl.hpp
+++ b/sycl/source/detail/device_impl.hpp
@@ -56,7 +56,8 @@ class device_impl {
   /// @return non-constant reference to PI device
   RT::PiDevice &getHandleRef() {
     if (MIsHostDevice)
-      throw invalid_object_error("This instance of device is a host instance");
+      throw invalid_object_error("This instance of device is a host instance",
+                                 PI_INVALID_DEVICE);
 
     return MDevice;
   }
@@ -68,7 +69,8 @@ class device_impl {
   /// @return constant reference to PI device
   const RT::PiDevice &getHandleRef() const {
     if (MIsHostDevice)
-      throw invalid_object_error("This instance of device is a host instance");
+      throw invalid_object_error("This instance of device is a host instance",
+                                 PI_INVALID_DEVICE);
 
     return MDevice;
   }
diff --git a/sycl/source/detail/device_info.cpp b/sycl/source/detail/device_info.cpp
index 1cdfee599fcc1..9135f0b1c9c92 100644
--- a/sycl/source/detail/device_info.cpp
+++ b/sycl/source/detail/device_info.cpp
@@ -35,7 +35,8 @@ device get_device_info<device, info::device::parent_device>::get(
       sizeof(result), &result, nullptr);
   if (result == nullptr)
     throw invalid_object_error(
-        "No parent for device because it is not a subdevice");
+        "No parent for device because it is not a subdevice",
+        PI_INVALID_DEVICE);
 
   return createSyclObjFromImpl<device>(
       std::make_shared<device_impl>(result, Plugin));
@@ -439,7 +440,8 @@ bool get_device_info_host<info::device::preferred_interop_user_sync>() {
 template <> device get_device_info_host<info::device::parent_device>() {
   // TODO: implement host device partitioning
   throw runtime_error(
-      "Partitioning to subdevices of the host device is not implemented yet");
+      "Partitioning to subdevices of the host device is not implemented yet",
+      PI_INVALID_DEVICE);
 }
 
 template <>
@@ -482,20 +484,23 @@ template <> cl_uint get_device_info_host<info::device::reference_count>() {
 
 template <> cl_uint get_device_info_host<info::device::max_num_sub_groups>() {
   // TODO update once subgroups are enabled
-  throw runtime_error("Sub-group feature is not supported on HOST device.");
+  throw runtime_error("Sub-group feature is not supported on HOST device.",
+                      PI_INVALID_DEVICE);
 }
 
 template <> vector_class<size_t>
 get_device_info_host<info::device::sub_group_sizes>() {
   // TODO update once subgroups are enabled
-  throw runtime_error("Sub-group feature is not supported on HOST device.");
+  throw runtime_error("Sub-group feature is not supported on HOST device.",
+                      PI_INVALID_DEVICE);
 }
 
 template <>
 bool get_device_info_host<
     info::device::sub_group_independent_forward_progress>() {
   // TODO update once subgroups are enabled
-  throw runtime_error("Sub-group feature is not supported on HOST device.");
+  throw runtime_error("Sub-group feature is not supported on HOST device.",
+                      PI_INVALID_DEVICE);
 }
 
 template <>
diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp
index 398e0cc7bdc9a..6f5c7920b4443 100644
--- a/sycl/source/detail/event_impl.cpp
+++ b/sycl/source/detail/event_impl.cpp
@@ -31,7 +31,8 @@ cl_event event_impl::get() const {
     return pi::cast<cl_event>(MEvent);
   }
   throw invalid_object_error(
-      "This instance of event doesn't support OpenCL interoperability.");
+      "This instance of event doesn't support OpenCL interoperability.",
+      PI_INVALID_EVENT);
 }
 
 event_impl::~event_impl() {
@@ -69,7 +70,8 @@ event_impl::event_impl(RT::PiEvent Event, const context &SyclContext)
   if (MContext->is_host()) {
     throw cl::sycl::invalid_parameter_error(
         "The syclContext must match the OpenCL context associated with the "
-        "clEvent.");
+        "clEvent.",
+        PI_INVALID_CONTEXT);
   }
 
   RT::PiContext TempContext;
@@ -78,7 +80,8 @@ event_impl::event_impl(RT::PiEvent Event, const context &SyclContext)
   if (MContext->getHandleRef() != TempContext) {
     throw cl::sycl::invalid_parameter_error(
         "The syclContext must match the OpenCL context associated with the "
-        "clEvent.");
+        "clEvent.",
+        PI_INVALID_CONTEXT);
   }
 
   getPlugin().call<PiApiKind::piEventRetain>(MEvent);
@@ -89,7 +92,7 @@ event_impl::event_impl(QueueImplPtr Queue) : MQueue(Queue) {
       Queue->has_property<property::queue::enable_profiling>()) {
     MHostProfilingInfo.reset(new HostProfilingInfo());
     if (!MHostProfilingInfo)
-      throw runtime_error("Out of host memory");
+      throw runtime_error("Out of host memory", PI_OUT_OF_HOST_MEMORY);
   }
 }
 
@@ -126,7 +129,8 @@ event_impl::get_profiling_info<info::event_profiling::command_submit>() const {
         this->getHandleRef(), this->getPlugin());
   }
   if (!MHostProfilingInfo)
-    throw invalid_object_error("Profiling info is not available.");
+    throw invalid_object_error("Profiling info is not available.",
+                               PI_PROFILING_INFO_NOT_AVAILABLE);
   return MHostProfilingInfo->getStartTime();
 }
 
@@ -138,7 +142,8 @@ event_impl::get_profiling_info<info::event_profiling::command_start>() const {
         this->getHandleRef(), this->getPlugin());
   }
   if (!MHostProfilingInfo)
-    throw invalid_object_error("Profiling info is not available.");
+    throw invalid_object_error("Profiling info is not available.",
+                               PI_PROFILING_INFO_NOT_AVAILABLE);
   return MHostProfilingInfo->getStartTime();
 }
 
@@ -150,7 +155,8 @@ event_impl::get_profiling_info<info::event_profiling::command_end>() const {
         this->getHandleRef(), this->getPlugin());
   }
   if (!MHostProfilingInfo)
-    throw invalid_object_error("Profiling info is not available.");
+    throw invalid_object_error("Profiling info is not available.",
+                               PI_PROFILING_INFO_NOT_AVAILABLE);
   return MHostProfilingInfo->getEndTime();
 }
 
diff --git a/sycl/source/detail/force_device.cpp b/sycl/source/detail/force_device.cpp
index f261f24480052..a88b129d86f8d 100644
--- a/sycl/source/detail/force_device.cpp
+++ b/sycl/source/detail/force_device.cpp
@@ -39,7 +39,8 @@ info::device_type get_forced_type() {
       return info::device_type::host;
     }
     throw cl::sycl::runtime_error("SYCL_DEVICE_TYPE is not recognized.  Must "
-                                  "be GPU, CPU, ACC or HOST.");
+                                  "be GPU, CPU, ACC or HOST.",
+                                  PI_INVALID_VALUE);
   }
   return info::device_type::all;
 }
diff --git a/sycl/source/detail/image_impl.cpp b/sycl/source/detail/image_impl.cpp
index 8102fec48ee3a..595df166fe997 100644
--- a/sycl/source/detail/image_impl.cpp
+++ b/sycl/source/detail/image_impl.cpp
@@ -312,14 +312,16 @@ bool image_impl<Dimensions>::checkImageDesc(const RT::PiMemImageDesc &Desc,
                                                              Desc.image_width))
     throw invalid_parameter_error(
         "For a 1D/2D image/image array, the width must be a Value >= 1 and "
-        "<= CL_DEVICE_IMAGE2D_MAX_WIDTH.");
+        "<= CL_DEVICE_IMAGE2D_MAX_WIDTH.",
+        PI_INVALID_VALUE);
 
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE3D) &&
       !checkImageValueRange<info::device::image3d_max_width>(Context,
                                                              Desc.image_width))
     throw invalid_parameter_error(
         "For a 3D image, the width must be a Value >= 1 and <= "
-        "CL_DEVICE_IMAGE3D_MAX_WIDTH");
+        "CL_DEVICE_IMAGE3D_MAX_WIDTH",
+        PI_INVALID_VALUE);
 
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE2D,
                PI_MEM_TYPE_IMAGE2D_ARRAY) &&
@@ -327,21 +329,24 @@ bool image_impl<Dimensions>::checkImageDesc(const RT::PiMemImageDesc &Desc,
           Context, Desc.image_height))
     throw invalid_parameter_error("For a 2D image or image array, the height "
                                   "must be a Value >= 1 and <= "
-                                  "CL_DEVICE_IMAGE2D_MAX_HEIGHT");
+                                  "CL_DEVICE_IMAGE2D_MAX_HEIGHT",
+                                  PI_INVALID_VALUE);
 
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE3D) &&
       !checkImageValueRange<info::device::image3d_max_height>(
           Context, Desc.image_height))
     throw invalid_parameter_error(
         "For a 3D image, the heightmust be a Value >= 1 and <= "
-        "CL_DEVICE_IMAGE3D_MAX_HEIGHT");
+        "CL_DEVICE_IMAGE3D_MAX_HEIGHT",
+        PI_INVALID_VALUE);
 
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE3D) &&
       !checkImageValueRange<info::device::image3d_max_depth>(Context,
                                                              Desc.image_depth))
     throw invalid_parameter_error(
         "For a 3D image, the depth must be a Value >= 1 and <= "
-        "CL_DEVICE_IMAGE3D_MAX_DEPTH");
+        "CL_DEVICE_IMAGE3D_MAX_DEPTH",
+        PI_INVALID_VALUE);
 
   if (checkAny(Desc.image_type, PI_MEM_TYPE_IMAGE1D_ARRAY,
                PI_MEM_TYPE_IMAGE2D_ARRAY) &&
@@ -349,27 +354,30 @@ bool image_impl<Dimensions>::checkImageDesc(const RT::PiMemImageDesc &Desc,
           Context, Desc.image_array_size))
     throw invalid_parameter_error(
         "For a 1D and 2D image array, the array_size must be a "
-        "Value >= 1 and <= "
-        "CL_DEVICE_IMAGE_MAX_ARRAY_SIZE.");
+        "Value >= 1 and <= CL_DEVICE_IMAGE_MAX_ARRAY_SIZE.",
+        PI_INVALID_VALUE);
 
   if ((nullptr == UserPtr) && (0 != Desc.image_row_pitch))
     throw invalid_parameter_error(
-        "The row_pitch must be 0 if host_ptr is nullptr.");
+        "The row_pitch must be 0 if host_ptr is nullptr.", PI_INVALID_VALUE);
 
   if ((nullptr == UserPtr) && (0 != Desc.image_slice_pitch))
     throw invalid_parameter_error(
-        "The slice_pitch must be 0 if host_ptr is nullptr.");
+        "The slice_pitch must be 0 if host_ptr is nullptr.", PI_INVALID_VALUE);
 
   if (0 != Desc.num_mip_levels)
-    throw invalid_parameter_error("The mip_levels must be 0.");
+    throw invalid_parameter_error("The mip_levels must be 0.",
+                                  PI_INVALID_VALUE);
 
   if (0 != Desc.num_samples)
-    throw invalid_parameter_error("The num_samples must be 0.");
+    throw invalid_parameter_error("The num_samples must be 0.",
+                                  PI_INVALID_VALUE);
 
   if (nullptr != Desc.buffer)
     throw invalid_parameter_error(
         "The buffer must be nullptr, because SYCL does not support "
-        "image creation from memory objects.");
+        "image creation from memory objects.",
+        PI_INVALID_VALUE);
 
   return true;
 }
@@ -387,7 +395,8 @@ bool image_impl<Dimensions>::checkImageFormat(
     throw invalid_parameter_error(
         "CL_INTENSITY or CL_LUMINANCE format can only be used if channel "
         "data type = CL_UNORM_INT8, CL_UNORM_INT16, CL_SNORM_INT8, "
-        "CL_SNORM_INT16, CL_HALF_FLOAT, or CL_FLOAT. ");
+        "CL_SNORM_INT16, CL_HALF_FLOAT, or CL_FLOAT.",
+        PI_INVALID_VALUE);
 
   if (checkAny(Format.image_channel_order, PI_IMAGE_CHANNEL_ORDER_RGB,
                PI_IMAGE_CHANNEL_ORDER_RGBx) &&
@@ -398,7 +407,8 @@ bool image_impl<Dimensions>::checkImageFormat(
     throw invalid_parameter_error(
         "CL_RGB or CL_RGBx	These formats can only be used if channel data "
         "type = CL_UNORM_SHORT_565, CL_UNORM_SHORT_555 or "
-        "CL_UNORM_INT_101010. ");
+        "CL_UNORM_INT_101010.",
+        PI_INVALID_VALUE);
 
   if (checkAny(Format.image_channel_order, PI_IMAGE_CHANNEL_ORDER_ARGB,
                PI_IMAGE_CHANNEL_ORDER_BGRA, PI_IMAGE_CHANNEL_ORDER_ABGR) &&
@@ -409,7 +419,8 @@ bool image_impl<Dimensions>::checkImageFormat(
     throw invalid_parameter_error(
         "CL_ARGB, CL_BGRA, CL_ABGR	These formats can only be used if "
         "channel data type = CL_UNORM_INT8, CL_SNORM_INT8, CL_SIGNED_INT8 "
-        "or CL_UNSIGNED_INT8.");
+        "or CL_UNSIGNED_INT8.",
+        PI_INVALID_VALUE);
 
   return true;
 }
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
index 107800a5cc9b7..9b66422978eb1 100644
--- a/sycl/source/detail/kernel_impl.cpp
+++ b/sycl/source/detail/kernel_impl.cpp
@@ -37,7 +37,8 @@ kernel_impl::kernel_impl(RT::PiKernel Kernel, ContextImplPtr ContextImpl,
       MKernel, PI_KERNEL_INFO_CONTEXT, sizeof(Context), &Context, nullptr);
   if (ContextImpl->getHandleRef() != Context)
     throw cl::sycl::invalid_parameter_error(
-        "Input context must be the same as the context of cl_kernel");
+        "Input context must be the same as the context of cl_kernel",
+        PI_INVALID_CONTEXT);
   getPlugin().call<PiApiKind::piKernelRetain>(MKernel);
 }
 
@@ -88,7 +89,8 @@ template <info::kernel_sub_group param>
 typename info::param_traits<info::kernel_sub_group, param>::return_type
 kernel_impl::get_sub_group_info(const device &Device) const {
   if (is_host()) {
-    throw runtime_error("Sub-group feature is not supported on HOST device.");
+    throw runtime_error("Sub-group feature is not supported on HOST device.",
+                        PI_INVALID_DEVICE);
   }
   return get_kernel_sub_group_info<
       typename info::param_traits<info::kernel_sub_group, param>::return_type,
@@ -103,7 +105,8 @@ kernel_impl::get_sub_group_info(
     typename info::param_traits<info::kernel_sub_group, param>::input_type
         Value) const {
   if (is_host()) {
-    throw runtime_error("Sub-group feature is not supported on HOST device.");
+    throw runtime_error("Sub-group feature is not supported on HOST device.",
+                        PI_INVALID_DEVICE);
   }
   return get_kernel_sub_group_info_with_input<
       typename info::param_traits<info::kernel_sub_group, param>::return_type,
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index 50e375e4fa831..d026b052aa089 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/sycl/source/detail/kernel_impl.hpp
@@ -73,7 +73,8 @@ class kernel_impl {
   /// @return a valid cl_kernel instance
   cl_kernel get() const {
     if (is_host())
-      throw invalid_object_error("This instance of kernel is a host instance");
+      throw invalid_object_error("This instance of kernel is a host instance",
+                                 PI_INVALID_KERNEL);
     getPlugin().call<PiApiKind::piKernelRetain>(MKernel);
     return pi::cast<cl_kernel>(MKernel);
   }
diff --git a/sycl/source/detail/kernel_info.cpp b/sycl/source/detail/kernel_info.cpp
index ff6fd5a64a3cb..bced197de0cb6 100644
--- a/sycl/source/detail/kernel_info.cpp
+++ b/sycl/source/detail/kernel_info.cpp
@@ -16,7 +16,8 @@ template <>
 cl::sycl::range<3>
 get_kernel_work_group_info_host<info::kernel_work_group::global_work_size>(
     const cl::sycl::device &Dev) {
-  throw invalid_object_error("This instance of kernel is a host instance");
+  throw invalid_object_error("This instance of kernel is a host instance",
+                             PI_INVALID_KERNEL);
 }
 
 template <>
diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp
index 577045ac4aaa9..54df3781580d4 100644
--- a/sycl/source/detail/memory_manager.cpp
+++ b/sycl/source/detail/memory_manager.cpp
@@ -200,7 +200,8 @@ void *MemoryManager::allocateMemSubBuffer(ContextImplPtr TargetContext,
   if (Error == PI_MISALIGNED_SUB_BUFFER_OFFSET)
     throw invalid_object_error(
         "Specified offset of the sub-buffer being constructed is not a "
-        "multiple of the memory base address alignment");
+        "multiple of the memory base address alignment",
+        PI_INVALID_VALUE);
   return NewMem;
 }
 
@@ -354,7 +355,8 @@ static void copyH2H(SYCLMemObjI *SYCLMemObj, char *SrcMem,
       (SrcOffset != id<3>{0, 0, 0} || DstOffset != id<3>{0, 0, 0} ||
        SrcSize != SrcAccessRange || DstSize != DstAccessRange)) {
     assert(!"Not supported configuration of memcpy requested");
-    throw runtime_error("Not supported configuration of memcpy requested");
+    throw runtime_error("Not supported configuration of memcpy requested",
+                        PI_INVALID_OPERATION);
   }
 
   DstOffset[0] *= DstElemSize;
@@ -427,7 +429,8 @@ void MemoryManager::fill(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
       return;
     }
     assert(!"Not supported configuration of fill requested");
-    throw runtime_error("Not supported configuration of fill requested");
+    throw runtime_error("Not supported configuration of fill requested",
+                        PI_INVALID_OPERATION);
   } else {
     Plugin.call<PiApiKind::piEnqueueMemImageFill>(
         Queue->getHandleRef(), pi::cast<RT::PiMem>(Mem), Pattern, &Offset[0],
@@ -443,7 +446,8 @@ void *MemoryManager::map(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
                          RT::PiEvent &OutEvent) {
   if (Queue->is_host()) {
     assert(!"Not supported configuration of map requested");
-    throw runtime_error("Not supported configuration of map requested");
+    throw runtime_error("Not supported configuration of map requested",
+                        PI_INVALID_OPERATION);
   }
 
   cl_map_flags Flags = 0;
diff --git a/sycl/source/detail/platform_impl.cpp b/sycl/source/detail/platform_impl.cpp
index c0d360f68a473..ac0274c0bd4c7 100644
--- a/sycl/source/detail/platform_impl.cpp
+++ b/sycl/source/detail/platform_impl.cpp
@@ -101,13 +101,15 @@ static std::vector<DevDescT> getWhiteListDesc() {
     }
 
     if (':' != *str)
-      throw sycl::runtime_error("Malformed device white list");
+      throw sycl::runtime_error("Malformed device white list",
+                                PI_INVALID_VALUE);
 
     // Skip ':'
     str += 1;
 
     if ('{' != *str || '{' != *(str + 1))
-      throw sycl::runtime_error("Malformed device white list");
+      throw sycl::runtime_error("Malformed device white list",
+                                PI_INVALID_VALUE);
 
     // Skip opening sequence "{{"
     str += 2;
@@ -119,7 +121,8 @@ static std::vector<DevDescT> getWhiteListDesc() {
       ++str;
 
     if ('\0' == *str)
-      throw sycl::runtime_error("Malformed device white list");
+      throw sycl::runtime_error("Malformed device white list",
+                                PI_INVALID_VALUE);
 
     *size = str - *valuePtr;
 
@@ -133,7 +136,8 @@ static std::vector<DevDescT> getWhiteListDesc() {
     if ('|' == *str)
       decDescs.emplace_back();
     else if (',' != *str)
-      throw sycl::runtime_error("Malformed device white list");
+      throw sycl::runtime_error("Malformed device white list",
+                                PI_INVALID_VALUE);
 
     ++str;
   }
diff --git a/sycl/source/detail/platform_impl.hpp b/sycl/source/detail/platform_impl.hpp
index 78ef151764d12..d47cf2ecb0abf 100644
--- a/sycl/source/detail/platform_impl.hpp
+++ b/sycl/source/detail/platform_impl.hpp
@@ -82,8 +82,8 @@ class platform_impl {
   /// @return an instance of OpenCL cl_platform_id.
   cl_platform_id get() const {
     if (is_host())
-      throw invalid_object_error(
-          "This instance of platform is a host instance");
+      throw invalid_object_error("This instance of platform is a host instance",
+                                 PI_INVALID_PLATFORM);
 
     return pi::cast<cl_platform_id>(MPlatform);
   }
@@ -97,8 +97,8 @@ class platform_impl {
   /// @return a raw plug-in platform handle.
   const RT::PiPlatform &getHandleRef() const {
     if (is_host())
-      throw invalid_object_error(
-          "This instance of platform is a host instance");
+      throw invalid_object_error("This instance of platform is a host instance",
+                                 PI_INVALID_PLATFORM);
 
     return MPlatform;
   }
diff --git a/sycl/source/detail/platform_util.cpp b/sycl/source/detail/platform_util.cpp
index 5ea2b5ae06c3d..068650692c8db 100644
--- a/sycl/source/detail/platform_util.cpp
+++ b/sycl/source/detail/platform_util.cpp
@@ -31,7 +31,8 @@ static void cpuid(uint32_t *CPUInfo, uint32_t Type, uint32_t SubType = 0) {
 
 uint32_t PlatformUtil::getMaxClockFrequency() {
   throw runtime_error(
-      "max_clock_frequency parameter is not supported for host device");
+      "max_clock_frequency parameter is not supported for host device",
+      PI_INVALID_DEVICE);
   uint32_t CPUInfo[4];
   string_class Buff(sizeof(CPUInfo) * 3 + 1, 0);
   size_t Offset = 0;
diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp
index 915c210393579..205c5da8c219c 100644
--- a/sycl/source/detail/program_impl.cpp
+++ b/sycl/source/detail/program_impl.cpp
@@ -33,7 +33,8 @@ program_impl::program_impl(
       MBuildOptions(LinkOptions) {
   // Verify arguments
   if (ProgramList.empty()) {
-    throw runtime_error("Non-empty vector of programs expected");
+    throw runtime_error("Non-empty vector of programs expected",
+                        PI_INVALID_VALUE);
   }
   MContext = ProgramList[0]->MContext;
   MDevices = ProgramList[0]->MDevices;
@@ -46,14 +47,16 @@ program_impl::program_impl(
     Prg->throw_if_state_is_not(program_state::compiled);
     if (Prg->MContext != MContext) {
       throw invalid_object_error(
-          "Not all programs are associated with the same context");
+          "Not all programs are associated with the same context",
+          PI_INVALID_PROGRAM);
     }
     if (!is_host()) {
       vector_class<device> PrgDevicesSorted =
           sort_devices_by_cl_device_id(Prg->MDevices);
       if (PrgDevicesSorted != DevicesSorted) {
         throw invalid_object_error(
-            "Not all programs are associated with the same devices");
+            "Not all programs are associated with the same devices",
+            PI_INVALID_PROGRAM);
       }
     }
   }
@@ -152,7 +155,8 @@ program_impl::~program_impl() {
 cl_program program_impl::get() const {
   throw_if_state_is(program_state::none);
   if (is_host()) {
-    throw invalid_object_error("This instance of program is a host instance");
+    throw invalid_object_error("This instance of program is a host instance",
+                               PI_INVALID_PROGRAM);
   }
   const detail::plugin &Plugin = getPlugin();
   Plugin.call<PiApiKind::piProgramRetain>(MProgram);
@@ -245,7 +249,8 @@ kernel program_impl::get_kernel(string_class KernelName,
   throw_if_state_is(program_state::none);
   if (is_host()) {
     if (IsCreatedFromSource)
-      throw invalid_object_error("This instance of program is a host instance");
+      throw invalid_object_error("This instance of program is a host instance",
+                                 PI_INVALID_PROGRAM);
 
     return createSyclObjFromImpl<kernel>(
         std::make_shared<kernel_impl>(MContext, PtrToSelf));
@@ -297,7 +302,8 @@ void program_impl::compile(const string_class &Options) {
   if (Err != PI_SUCCESS) {
     throw compile_program_error(
         "Program compilation error:\n" +
-        ProgramManager::getProgramBuildLog(MProgram, MContext));
+            ProgramManager::getProgramBuildLog(MProgram, MContext),
+        Err);
   }
   MCompileOptions = Options;
   MBuildOptions = Options;
@@ -314,7 +320,8 @@ void program_impl::build(const string_class &Options) {
   if (Err != PI_SUCCESS) {
     throw compile_program_error(
         "Program build error:\n" +
-        ProgramManager::getProgramBuildLog(MProgram, MContext));
+            ProgramManager::getProgramBuildLog(MProgram, MContext),
+        Err);
   }
   MBuildOptions = Options;
 }
@@ -359,7 +366,8 @@ RT::PiKernel program_impl::get_pi_kernel(const string_class &KernelName) const {
         MProgram, KernelName.c_str(), &Kernel);
     if (Err == PI_RESULT_INVALID_KERNEL_NAME) {
       throw invalid_object_error(
-          "This instance of program does not contain the kernel requested");
+          "This instance of program does not contain the kernel requested",
+          Err);
     }
     Plugin.checkPiResult(Err);
   }
@@ -379,13 +387,13 @@ program_impl::sort_devices_by_cl_device_id(vector_class<device> Devices) {
 
 void program_impl::throw_if_state_is(program_state State) const {
   if (MState == State) {
-    throw invalid_object_error("Invalid program state");
+    throw invalid_object_error("Invalid program state", PI_INVALID_PROGRAM);
   }
 }
 
 void program_impl::throw_if_state_is_not(program_state State) const {
   if (MState != State) {
-    throw invalid_object_error("Invalid program state");
+    throw invalid_object_error("Invalid program state", PI_INVALID_PROGRAM);
   }
 }
 
@@ -400,7 +408,8 @@ void program_impl::create_pi_program_with_kernel_name(
 template <>
 cl_uint program_impl::get_info<info::program::reference_count>() const {
   if (is_host()) {
-    throw invalid_object_error("This instance of program is a host instance");
+    throw invalid_object_error("This instance of program is a host instance",
+                               PI_INVALID_PROGRAM);
   }
   cl_uint Result;
   const detail::plugin &Plugin = getPlugin();
diff --git a/sycl/source/detail/program_impl.hpp b/sycl/source/detail/program_impl.hpp
index 9580d2fb20994..2727a8e44d500 100644
--- a/sycl/source/detail/program_impl.hpp
+++ b/sycl/source/detail/program_impl.hpp
@@ -298,7 +298,8 @@ class program_impl {
     for (const auto &Device : Devices) {
       if (!Device.get_info<param>()) {
         throw feature_not_supported(
-            "Online compilation is not supported by this device");
+            "Online compilation is not supported by this device",
+            PI_COMPILER_NOT_AVAILABLE);
       }
     }
   }
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index e2959329ce441..6b369bd0f9bb7 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -306,10 +306,12 @@ RT::PiProgram ProgramManager::createPIProgram(const DeviceImage &Img,
 
   // perform minimal sanity checks on the device image and the descriptor
   if (Img.BinaryEnd < Img.BinaryStart) {
-    throw runtime_error("Malformed device program image descriptor");
+    throw runtime_error("Malformed device program image descriptor",
+                        PI_INVALID_VALUE);
   }
   if (Img.BinaryEnd == Img.BinaryStart) {
-    throw runtime_error("Invalid device program image: size is zero");
+    throw runtime_error("Invalid device program image: size is zero",
+                        PI_INVALID_VALUE);
   }
   size_t ImgSize = static_cast<size_t>(Img.BinaryEnd - Img.BinaryStart);
 
@@ -325,7 +327,8 @@ RT::PiProgram ProgramManager::createPIProgram(const DeviceImage &Img,
 
   if (!isDeviceBinaryTypeSupported(Context, Format))
     throw feature_not_supported(
-        "Online compilation is not supported in this context");
+        "Online compilation is not supported in this context",
+        PI_INVALID_OPERATION);
 
   // Load the image
   const ContextImplPtr Ctx = getSyclObjImpl(Context);
@@ -503,7 +506,8 @@ static const char* getDeviceLibFilename(DeviceLibExt Extension) {
   case cl_intel_devicelib_complex_fp64:
     return "libsycl-fallback-complex-fp64.spv";
   }
-  throw compile_program_error("Unhandled (new?) device library extension");
+  throw compile_program_error("Unhandled (new?) device library extension",
+                              PI_INVALID_OPERATION);
 }
 
 static const char* getDeviceLibExtensionStr(DeviceLibExt Extension) {
@@ -519,7 +523,8 @@ static const char* getDeviceLibExtensionStr(DeviceLibExt Extension) {
   case cl_intel_devicelib_complex_fp64:
     return "cl_intel_devicelib_complex_fp64";
   }
-  throw compile_program_error("Unhandled (new?) device library extension");
+  throw compile_program_error("Unhandled (new?) device library extension",
+                              PI_INVALID_OPERATION);
 }
 
 static RT::PiProgram loadDeviceLibFallback(
@@ -540,7 +545,8 @@ static RT::PiProgram loadDeviceLibFallback(
 
   if (!loadDeviceLib(Context, LibFileName, LibProg)) {
     CachedLibPrograms.erase(LibProgIt);
-    throw compile_program_error(std::string("Failed to load ") + LibFileName);
+    throw compile_program_error(std::string("Failed to load ") + LibFileName,
+                                PI_INVALID_VALUE);
   }
 
   const detail::plugin &Plugin = Context->getPlugin();
@@ -556,7 +562,7 @@ static RT::PiProgram loadDeviceLibFallback(
   if (Error != PI_SUCCESS) {
     CachedLibPrograms.erase(LibProgIt);
     throw compile_program_error(
-        ProgramManager::getProgramBuildLog(LibProg, Context));
+        ProgramManager::getProgramBuildLog(LibProg, Context), Error);
   }
 
   return LibProg;
@@ -580,7 +586,8 @@ ProgramManager::ProgramManager() {
 
     if (!File.is_open())
       throw runtime_error(std::string("Can't open file specified via ") +
-                          UseSpvEnv + ": " + SpvFile);
+                              UseSpvEnv + ": " + SpvFile,
+                          PI_INVALID_VALUE);
     File.seekg(0, std::ios::end);
     size_t Size = File.tellg();
     std::unique_ptr<unsigned char[]> Data(new unsigned char[Size]);
@@ -589,7 +596,8 @@ ProgramManager::ProgramManager() {
     File.close();
     if (!File.good())
       throw runtime_error(std::string("read from ") + SpvFile +
-                          std::string(" failed"));
+                              std::string(" failed"),
+                          PI_INVALID_VALUE);
 
     std::unique_ptr<DeviceImage, ImageDeleter> ImgPtr(new DeviceImage(),
                                                       ImageDeleter());
@@ -762,7 +770,8 @@ ProgramManager::build(ProgramPtr Program, const ContextImplPtr Context,
         Program.get(), Devices.size(), Devices.data(), Opts.c_str(), nullptr,
         nullptr);
     if (Error != PI_SUCCESS)
-      throw compile_program_error(getProgramBuildLog(Program.get(), Context));
+      throw compile_program_error(getProgramBuildLog(Program.get(), Context),
+                                  Error);
     return Program;
   }
 
@@ -784,7 +793,8 @@ ProgramManager::build(ProgramPtr Program, const ContextImplPtr Context,
     if (LinkedProg) {
       // A non-trivial error occurred during linkage: get a build log, release
       // an incomplete (but valid) LinkedProg, and throw.
-      throw compile_program_error(getProgramBuildLog(LinkedProg, Context));
+      throw compile_program_error(getProgramBuildLog(LinkedProg, Context),
+                                  Error);
     }
     Plugin.checkPiResult(Error);
   }
@@ -898,7 +908,8 @@ ProgramManager::getKernelSetId(OSModuleHandle M,
   if (ModuleKSIdIt != m_OSModuleKernelSets.end())
     return ModuleKSIdIt->second;
 
-  throw runtime_error("No kernel named " + KernelName + " was found");
+  throw runtime_error("No kernel named " + KernelName + " was found",
+                      PI_RESULT_INVALID_KERNEL_NAME);
 }
 
 RT::PiDeviceBinaryType ProgramManager::getFormat(const DeviceImage &Img) const {
@@ -948,7 +959,8 @@ void ProgramManager::dumpImage(const DeviceImage &Img, KernelSetId KSId) const {
   std::ofstream F(Fname, std::ios::binary);
 
   if (!F.is_open()) {
-    throw runtime_error(std::string("Can not write ") + Fname);
+    throw runtime_error(std::string("Can not write ") + Fname,
+                        PI_INVALID_VALUE);
   }
   size_t ImgSize = static_cast<size_t>(Img.BinaryEnd - Img.BinaryStart);
   F.write(reinterpret_cast<const char *>(Img.BinaryStart), ImgSize);
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
index 2d1fd58e8489a..701c1f6941694 100644
--- a/sycl/source/detail/queue_impl.hpp
+++ b/sycl/source/detail/queue_impl.hpp
@@ -75,7 +75,8 @@ class queue_impl {
     if (!Context->hasDevice(Device))
       throw cl::sycl::invalid_parameter_error(
           "Queue cannot be constructed with the given context and device "
-          "as the context does not contain the given device.");
+          "as the context does not contain the given device.",
+          PI_INVALID_DEVICE);
   }
 
   /// Constructs a SYCL queue from plugin interoperability handle.
@@ -117,7 +118,8 @@ class queue_impl {
       return pi::cast<cl_command_queue>(MCommandQueue);
     }
     throw invalid_object_error(
-        "This instance of queue doesn't support OpenCL interoperability");
+        "This instance of queue doesn't support OpenCL interoperability",
+        PI_INVALID_QUEUE);
   }
 
   /// @return an associated SYCL context.
diff --git a/sycl/source/detail/sampler_impl.cpp b/sycl/source/detail/sampler_impl.cpp
index bad995ad752bc..1e875839060e1 100644
--- a/sycl/source/detail/sampler_impl.cpp
+++ b/sycl/source/detail/sampler_impl.cpp
@@ -65,7 +65,8 @@ RT::PiSampler sampler_impl::getOrCreateSampler(const context &Context) {
       getSyclObjImpl(Context)->getHandleRef(), sprops, &resultSampler);
 
   if (errcode_ret == PI_INVALID_OPERATION)
-    throw feature_not_supported("Images are not supported by this device.");
+    throw feature_not_supported("Images are not supported by this device.",
+                                errcode_ret);
 
   Plugin.checkPiResult(errcode_ret);
   m_contextToSampler[Context] = resultSampler;
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index 3b97de5673a92..6343ddecc670a 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -507,7 +507,7 @@ AllocaCommandBase *ExecCGCommand::getAllocaForReq(Requirement *Req) {
     if (Dep.MDepRequirement == Req)
       return Dep.MAllocaCmd;
   }
-  throw runtime_error("Alloca for command not found");
+  throw runtime_error("Alloca for command not found", PI_INVALID_OPERATION);
 }
 
 void ExecCGCommand::flushStreams() {
@@ -765,7 +765,8 @@ cl_int ExecCGCommand::enqueueImp() {
 
   case CG::CGTYPE::UPDATE_HOST: {
     assert(!"Update host should be handled by the Scheduler.");
-    throw runtime_error("Update host should be handled by the Scheduler.");
+    throw runtime_error("Update host should be handled by the Scheduler.",
+                        PI_INVALID_OPERATION);
   }
   case CG::CGTYPE::COPY_ACC_TO_PTR: {
     CGCopy *Copy = (CGCopy *)MCommandGroup.get();
@@ -1043,7 +1044,7 @@ cl_int ExecCGCommand::enqueueImp() {
   }
   case CG::CGTYPE::NONE:
   default:
-    throw runtime_error("CG type not implemented.");
+    throw runtime_error("CG type not implemented.", PI_INVALID_OPERATION);
   }
 }
 
diff --git a/sycl/source/detail/scheduler/graph_builder.cpp b/sycl/source/detail/scheduler/graph_builder.cpp
index d112de043a99a..62db768e40dfc 100644
--- a/sycl/source/detail/scheduler/graph_builder.cpp
+++ b/sycl/source/detail/scheduler/graph_builder.cpp
@@ -230,7 +230,7 @@ Command *Scheduler::GraphBuilder::insertMemoryMove(MemObjRecord *Record,
 
   AllocaCommandBase *AllocaCmdDst = getOrCreateAllocaForReq(Record, Req, Queue);
   if (!AllocaCmdDst)
-    throw runtime_error("Out of host memory");
+    throw runtime_error("Out of host memory", PI_OUT_OF_HOST_MEMORY);
 
   std::set<Command *> Deps =
       findDepsForReq(Record, Req, Queue->getContextImplPtr());
@@ -264,7 +264,7 @@ Command *Scheduler::GraphBuilder::insertMemoryMove(MemObjRecord *Record,
     AllocaCmdSrc = (Record->MAllocaCommands.end() != It) ? *It : nullptr;
   }
   if (!AllocaCmdSrc)
-    throw runtime_error("Cannot find buffer allocation");
+    throw runtime_error("Cannot find buffer allocation", PI_INVALID_VALUE);
   // Get parent allocation of sub buffer to perform full copy of whole buffer
   if (IsSuitableSubReq(Req)) {
     if (AllocaCmdSrc->getType() == Command::CommandType::ALLOCA_SUB_BUF)
@@ -322,7 +322,7 @@ Command *Scheduler::GraphBuilder::addCopyBack(Requirement *Req) {
       SrcAllocaCmd->getQueue(), std::move(HostQueue)));
 
   if (!MemCpyCmdUniquePtr)
-    throw runtime_error("Out of host memory");
+    throw runtime_error("Out of host memory", PI_OUT_OF_HOST_MEMORY);
 
   MemCpyCommandHost *MemCpyCmd = MemCpyCmdUniquePtr.release();
   for (Command *Dep : Deps) {
@@ -591,7 +591,7 @@ Scheduler::GraphBuilder::addCG(std::unique_ptr<detail::CG> CommandGroup,
   std::unique_ptr<ExecCGCommand> NewCmd(
       new ExecCGCommand(std::move(CommandGroup), Queue));
   if (!NewCmd)
-    throw runtime_error("Out of host memory");
+    throw runtime_error("Out of host memory", PI_OUT_OF_HOST_MEMORY);
 
   if (MPrintOptionsArray[BeforeAddCG])
     printGraphAsDot("before_addCG");
diff --git a/sycl/source/detail/scheduler/graph_processor.cpp b/sycl/source/detail/scheduler/graph_processor.cpp
index 903f4ac43294a..95f393afaff8f 100644
--- a/sycl/source/detail/scheduler/graph_processor.cpp
+++ b/sycl/source/detail/scheduler/graph_processor.cpp
@@ -43,7 +43,7 @@ void Scheduler::GraphProcessor::waitForEvent(EventImplPtr Event) {
   bool Enqueued = enqueueCommand(Cmd, Res, BLOCKING);
   if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
     // TODO: Reschedule commands.
-    throw runtime_error("Enqueue process failed.");
+    throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
 
   RT::PiEvent &CLEvent = Cmd->getEvent()->getHandleRef();
   if (CLEvent) {
diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp
index 4b9ce7e017f29..862507c1839cc 100644
--- a/sycl/source/detail/scheduler/scheduler.cpp
+++ b/sycl/source/detail/scheduler/scheduler.cpp
@@ -20,8 +20,8 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
-EventImplPtr addHostAccessorToSchedulerInstance(Requirement *Req, 
-                                               const bool destructor) {
+EventImplPtr addHostAccessorToSchedulerInstance(Requirement *Req,
+                                                const bool destructor) {
   return cl::sycl::detail::Scheduler::getInstance().
                                               addHostAccessor(Req, destructor);
 }
@@ -31,14 +31,14 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record) {
     EnqueueResultT Res;
     bool Enqueued = GraphProcessor::enqueueCommand(Cmd, Res);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.");
+      throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
     GraphProcessor::waitForEvent(Cmd->getEvent());
   }
   for (Command *Cmd : Record->MWriteLeaves) {
     EnqueueResultT Res;
     bool Enqueued = GraphProcessor::enqueueCommand(Cmd, Res);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.");
+      throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
     GraphProcessor::waitForEvent(Cmd->getEvent());
   }
   for (AllocaCommandBase *AllocaCmd : Record->MAllocaCommands) {
@@ -46,7 +46,7 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record) {
     EnqueueResultT Res;
     bool Enqueued = GraphProcessor::enqueueCommand(ReleaseCmd, Res);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.");
+      throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
     GraphProcessor::waitForEvent(ReleaseCmd->getEvent());
   }
 }
@@ -71,7 +71,7 @@ EventImplPtr Scheduler::addCG(std::unique_ptr<detail::CG> CommandGroup,
     EnqueueResultT Res;
     bool Enqueued = GraphProcessor::enqueueCommand(NewCmd, Res);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.");
+      throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
   }
 
   if (IsKernel)
@@ -92,7 +92,7 @@ EventImplPtr Scheduler::addCopyBack(Requirement *Req) {
     EnqueueResultT Res;
     bool Enqueued = GraphProcessor::enqueueCommand(NewCmd, Res);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.");
+      throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
   } catch (...) {
     NewCmd->getQueue()->reportAsyncException(std::current_exception());
   }
@@ -145,7 +145,7 @@ void Scheduler::removeMemoryObject(detail::SYCLMemObjI *MemObj) {
   MGraphBuilder.removeRecordForMemObj(MemObj);
 }
 
-EventImplPtr Scheduler::addHostAccessor(Requirement *Req, 
+EventImplPtr Scheduler::addHostAccessor(Requirement *Req,
                                         const bool destructor) {
   std::lock_guard<std::mutex> lock(MGraphLock);
 
@@ -156,7 +156,7 @@ EventImplPtr Scheduler::addHostAccessor(Requirement *Req,
   EnqueueResultT Res;
   bool Enqueued = GraphProcessor::enqueueCommand(NewCmd, Res);
   if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-    throw runtime_error("Enqueue process failed.");
+    throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
   return NewCmd->getEvent();
 }
 
@@ -168,7 +168,7 @@ void Scheduler::releaseHostAccessor(Requirement *Req) {
       EnqueueResultT Res;
       bool Enqueued = GraphProcessor::enqueueCommand(Cmd, Res);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.");
+        throw runtime_error("Enqueue process failed.", PI_INVALID_OPERATION);
     }
   };
   EnqueueLeaves(Record->MReadLeaves);
diff --git a/sycl/source/detail/sycl_mem_obj_t.cpp b/sycl/source/detail/sycl_mem_obj_t.cpp
index 142803dd33290..8f5b3e2c1ade2 100644
--- a/sycl/source/detail/sycl_mem_obj_t.cpp
+++ b/sycl/source/detail/sycl_mem_obj_t.cpp
@@ -29,17 +29,19 @@ SYCLMemObjT::SYCLMemObjT(cl_mem MemObject, const context &SyclContext,
   if (MInteropContext->is_host())
     throw cl::sycl::invalid_parameter_error(
         "Creation of interoperability memory object using host context is "
-        "not allowed");
+        "not allowed",
+        PI_INVALID_CONTEXT);
 
   RT::PiMem Mem = pi::cast<RT::PiMem>(MInteropMemObject);
   RT::PiContext Context = nullptr;
   const plugin &Plugin = getPlugin();
   Plugin.call<PiApiKind::piMemGetInfo>(Mem, CL_MEM_CONTEXT, sizeof(Context),
-                                         &Context, nullptr);
+                                       &Context, nullptr);
 
   if (MInteropContext->getHandleRef() != Context)
     throw cl::sycl::invalid_parameter_error(
-        "Input context must be the same as the context of cl_mem");
+        "Input context must be the same as the context of cl_mem",
+        PI_INVALID_CONTEXT);
   Plugin.call<PiApiKind::piMemRetain>(Mem);
 }
 
diff --git a/sycl/source/detail/usm/usm_impl.cpp b/sycl/source/detail/usm/usm_impl.cpp
index e6cb56c936cd6..b81529e553cfa 100644
--- a/sycl/source/detail/usm/usm_impl.cpp
+++ b/sycl/source/detail/usm/usm_impl.cpp
@@ -293,7 +293,7 @@ alloc get_pointer_type(const void *Ptr, const context &Ctxt) {
 device get_pointer_device(const void *Ptr, const context &Ctxt) {
   // Check if ptr is a valid USM pointer
   if (get_pointer_type(Ptr, Ctxt) == alloc::unknown)
-    throw runtime_error("Ptr not a valid USM allocation!");
+    throw runtime_error("Ptr not a valid USM allocation!", PI_INVALID_VALUE);
 
   // Just return the host device in the host context
   if (Ctxt.is_host())
@@ -305,7 +305,7 @@ device get_pointer_device(const void *Ptr, const context &Ctxt) {
   if (get_pointer_type(Ptr, Ctxt) == alloc::host) {
     auto Devs = CtxImpl->getDevices();
     if (Devs.size() == 0)
-      throw runtime_error("No devices in passed context!");
+      throw runtime_error("No devices in passed context!", PI_INVALID_VALUE);
 
     // Just return the first device in the context
     return Devs[0];
@@ -325,7 +325,8 @@ device get_pointer_device(const void *Ptr, const context &Ctxt) {
       return Dev;
   }
 
-  throw runtime_error("Cannot find device associated with USM allocation!");
+  throw runtime_error("Cannot find device associated with USM allocation!",
+                      PI_INVALID_OPERATION);
 }
 
 } // namespace sycl
diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp
index 6da02b72eeeda..7f85c46b4ad50 100644
--- a/sycl/source/device.cpp
+++ b/sycl/source/device.cpp
@@ -19,7 +19,8 @@ void force_type(info::device_type &t, const info::device_type &ft) {
   if (t == info::device_type::all) {
     t = ft;
   } else if (ft != info::device_type::all && t != ft) {
-    throw cl::sycl::invalid_parameter_error("No device of forced type.");
+    throw cl::sycl::invalid_parameter_error("No device of forced type.",
+                                            PI_INVALID_OPERATION);
   }
 }
 } // namespace detail
diff --git a/sycl/source/device_selector.cpp b/sycl/source/device_selector.cpp
index c08530b2d1163..1b6c2d7bef7ad 100644
--- a/sycl/source/device_selector.cpp
+++ b/sycl/source/device_selector.cpp
@@ -27,7 +27,8 @@ device device_selector::select_device() const {
   if (res != nullptr)
     return *res;
 
-  throw cl::sycl::runtime_error("No device of requested type available.");
+  throw cl::sycl::runtime_error("No device of requested type available.",
+                                PI_DEVICE_NOT_FOUND);
 }
 
 int default_selector::operator()(const device &dev) const {
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index 973d7262da6df..9d49f0c3e8efa 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -79,9 +79,11 @@ event handler::finalize() {
     break;
   case detail::CG::NONE:
     throw runtime_error("Command group submitted without a kernel or a "
-                        "explicit memory operation.");
+                        "explicit memory operation.",
+                        PI_INVALID_OPERATION);
   default:
-    throw runtime_error("Unhandled type of command group");
+    throw runtime_error("Unhandled type of command group",
+                        PI_INVALID_OPERATION);
   }
 
   detail::EventImplPtr Event = detail::Scheduler::getInstance().addCG(
@@ -174,7 +176,7 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
     case access::target::host_image:
     case access::target::host_buffer: {
       throw cl::sycl::invalid_parameter_error(
-          "Unsupported accessor target case.");
+          "Unsupported accessor target case.", PI_INVALID_OPERATION);
       break;
     }
     }
diff --git a/sycl/source/ordered_queue.cpp b/sycl/source/ordered_queue.cpp
index 228bd27a3f890..4abf0923fef49 100644
--- a/sycl/source/ordered_queue.cpp
+++ b/sycl/source/ordered_queue.cpp
@@ -59,7 +59,8 @@ ordered_queue::ordered_queue(cl_command_queue clQueue,
       &reportedProps, nullptr);
   if (reportedProps & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE)
     throw runtime_error(
-        "Failed to build a sycl ordered queue from a cl OOO queue.");
+        "Failed to build a sycl ordered queue from a cl OOO queue.",
+        PI_INVALID_OPERATION);
 
   impl = std::make_shared<detail::queue_impl>(
       m_CommandQueue, detail::getSyclObjImpl(syclContext), asyncHandler);
diff --git a/sycl/test/sub_group/attributes.cpp b/sycl/test/sub_group/attributes.cpp
index e032de24f30db..027327324329e 100644
--- a/sycl/test/sub_group/attributes.cpp
+++ b/sycl/test/sub_group/attributes.cpp
@@ -101,7 +101,8 @@ int main() {
         submit<KernelFunctor1>(Queue);
         break;
       default:
-        throw feature_not_supported("sub-group size is not supported");
+        throw feature_not_supported("sub-group size is not supported",
+                                    PI_INVALID_OPERATION);
       }
 
       auto Kernel = TheKernel[0];
diff --git a/sycl/test/sub_group/common_ocl.cpp b/sycl/test/sub_group/common_ocl.cpp
index 9a8e4afe7cd3f..ca033fe955bf3 100644
--- a/sycl/test/sub_group/common_ocl.cpp
+++ b/sycl/test/sub_group/common_ocl.cpp
@@ -39,7 +39,7 @@ void check(queue &Queue, const int G, const int L, const char *SpvFile) {
     std::ifstream File(SpvFile, std::ios::binary);
     if (!File.is_open()) {
       std::cerr << std::strerror(errno);
-      throw compile_program_error("Cannot open SPIRV file\n");
+      throw compile_program_error("Cannot open SPIRV file\n", PI_INVALID_VALUE);
     }
     File.seekg(0, std::ios::end);
     vector_class<char> Spv(File.tellg());