diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 320b828336c99..e01ba11035608 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -843,8 +843,9 @@ class image_accessor
 
   sycl::vec<int, Dimensions> getRangeInternal() const {
     // TODO: Implement for host.
-    throw runtime_error("image::getRangeInternal() is not implemented for host",
-                        PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(
+        make_error_code(errc::feature_not_supported),
+        "image::getRangeInternal() is not implemented for host");
     return sycl::vec<int, Dimensions>{1};
   }
 
diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp
index 8ed58d3625a7c..a33cc4367f8fc 100644
--- a/sycl/include/sycl/backend.hpp
+++ b/sycl/include/sycl/backend.hpp
@@ -114,10 +114,8 @@ auto get_native_buffer(const buffer<DataT, Dimensions, AllocatorT, void> &Obj)
   // No check for backend mismatch because buffer can be allocated on different
   // backends
   if (BackendName == backend::ext_oneapi_level_zero)
-    throw sycl::runtime_error(
-        errc::feature_not_supported,
-        "Buffer interop is not supported by level zero yet",
-        PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Buffer interop is not supported by level zero yet");
   return Obj.template getNative<BackendName>();
 }
 #endif
@@ -126,10 +124,9 @@ auto get_native_buffer(const buffer<DataT, Dimensions, AllocatorT, void> &Obj)
 template <backend BackendName, class SyclObjectT>
 auto get_native(const SyclObjectT &Obj)
     -> backend_return_t<BackendName, SyclObjectT> {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != BackendName) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::backend_mismatch),
+                          "Backends mismatch");
   }
   return reinterpret_cast<backend_return_t<BackendName, SyclObjectT>>(
       Obj.getNative());
@@ -137,10 +134,9 @@ auto get_native(const SyclObjectT &Obj)
 
 template <backend BackendName>
 auto get_native(const queue &Obj) -> backend_return_t<BackendName, queue> {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != BackendName) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::backend_mismatch),
+                          "Backends mismatch");
   }
   int32_t IsImmCmdList;
   pi_native_handle Handle = Obj.getNative(IsImmCmdList);
@@ -160,10 +156,9 @@ auto get_native(const queue &Obj) -> backend_return_t<BackendName, queue> {
 template <backend BackendName, bundle_state State>
 auto get_native(const kernel_bundle<State> &Obj)
     -> backend_return_t<BackendName, kernel_bundle<State>> {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != BackendName) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::backend_mismatch),
+                          "Backends mismatch");
   }
   return Obj.template getNative<BackendName>();
 }
@@ -179,10 +174,9 @@ auto get_native(const buffer<DataT, Dimensions, AllocatorT> &Obj)
 template <>
 inline backend_return_t<backend::opencl, event>
 get_native<backend::opencl, event>(const event &Obj) {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != backend::opencl) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::backend_mismatch),
+                          "Backends mismatch");
   }
   backend_return_t<backend::opencl, event> ReturnValue;
   for (auto const &element : Obj.getNativeVector()) {
@@ -199,10 +193,9 @@ get_native<backend::opencl, event>(const event &Obj) {
 template <>
 inline backend_return_t<backend::ext_oneapi_cuda, device>
 get_native<backend::ext_oneapi_cuda, device>(const device &Obj) {
-  // TODO use SYCL 2020 exception when implemented
   if (Obj.get_backend() != backend::ext_oneapi_cuda) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw sycl::exception(make_error_code(errc::backend_mismatch),
+                          "Backends mismatch");
   }
   // CUDA uses a 32-bit int instead of an opaque pointer like other backends,
   // so we need a specialization with static_cast instead of reinterpret_cast.
diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp
index 61b3684daf65f..f244ba691d719 100644
--- a/sycl/include/sycl/group.hpp
+++ b/sycl/include/sycl/group.hpp
@@ -11,12 +11,12 @@
 #include <CL/__spirv/spirv_ops.hpp>
 #include <CL/__spirv/spirv_types.hpp>
 #include <CL/__spirv/spirv_vars.hpp>
-#include <stdexcept>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/spirv.hpp>
 #include <sycl/device_event.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/h_item.hpp>
 #include <sycl/id.hpp>
 #include <sycl/memory_enums.hpp>
@@ -24,6 +24,8 @@
 #include <sycl/range.hpp>
 #include <type_traits>
 
+#include <stdexcept>
+
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 namespace detail {
@@ -127,8 +129,8 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv::initLocalInvocationId<Dimensions, id<Dimensions>>();
 #else
-    throw runtime_error("get_local_id() is not implemented on host",
-                        PI_ERROR_INVALID_DEVICE);
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "get_local_id() is not implemented on host");
 #endif
   }
 
diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp
index 08e40f0dab518..dfc0a4c407c05 100644
--- a/sycl/include/sycl/group_algorithm.hpp
+++ b/sycl/include/sycl/group_algorithm.hpp
@@ -15,6 +15,7 @@
 #include <sycl/builtins.hpp>
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp>
 #include <sycl/ext/oneapi/functional.hpp>
 #include <sycl/functional.hpp>
@@ -196,8 +197,8 @@ Function for_each(Group g, Ptr first, Ptr last, Function f) {
   (void)first;
   (void)last;
   (void)f;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 } // namespace detail
@@ -233,8 +234,8 @@ reduce_over_group(Group g, T x, BinaryOperation binary_op) {
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -256,8 +257,8 @@ reduce_over_group(Group g, T x, BinaryOperation binary_op) {
   (void)g;
   (void)x;
   (void)binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -298,8 +299,8 @@ reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
   return binary_op(init, reduce_over_group(g, T(x), binary_op));
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -325,8 +326,8 @@ reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
   return result;
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -354,8 +355,8 @@ joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
 #else
   (void)g;
   (void)last;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -377,8 +378,8 @@ joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
   (void)first;
   (void)last;
   (void)binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -397,8 +398,8 @@ any_of_group(Group g, bool pred) {
 #else
   (void)g;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -423,8 +424,8 @@ joint_any_of(Group g, Ptr first, Ptr last, Predicate pred) {
   (void)first;
   (void)last;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -443,8 +444,8 @@ all_of_group(Group g, bool pred) {
 #else
   (void)g;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -469,8 +470,8 @@ joint_all_of(Group g, Ptr first, Ptr last, Predicate pred) {
   (void)first;
   (void)last;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -489,8 +490,8 @@ none_of_group(Group g, bool pred) {
 #else
   (void)g;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -512,8 +513,8 @@ joint_none_of(Group g, Ptr first, Ptr last, Predicate pred) {
   (void)first;
   (void)last;
   (void)pred;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -531,8 +532,8 @@ shift_group_left(Group, T x, typename Group::linear_id_type delta = 1) {
 #else
   (void)x;
   (void)delta;
-  throw runtime_error("Sub-groups are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
 #endif
 }
 
@@ -550,8 +551,8 @@ shift_group_right(Group, T x, typename Group::linear_id_type delta = 1) {
 #else
   (void)x;
   (void)delta;
-  throw runtime_error("Sub-groups are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
 #endif
 }
 
@@ -569,8 +570,8 @@ permute_group_by_xor(Group, T x, typename Group::linear_id_type mask) {
 #else
   (void)x;
   (void)mask;
-  throw runtime_error("Sub-groups are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
 #endif
 }
 
@@ -588,8 +589,8 @@ select_from_group(Group, T x, typename Group::id_type local_id) {
 #else
   (void)x;
   (void)local_id;
-  throw runtime_error("Sub-groups are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
 #endif
 }
 
@@ -615,8 +616,8 @@ group_broadcast(Group g, T x, typename Group::id_type local_id) {
   (void)g;
   (void)x;
   (void)local_id;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -634,8 +635,8 @@ group_broadcast(Group g, T x, typename Group::linear_id_type linear_local_id) {
   (void)g;
   (void)x;
   (void)linear_local_id;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -650,8 +651,8 @@ group_broadcast(Group g, T x) {
 #else
   (void)g;
   (void)x;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -681,8 +682,8 @@ exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -704,8 +705,8 @@ exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
   (void)g;
   (void)x;
   (void)binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -773,8 +774,8 @@ exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
   return scan;
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -824,8 +825,8 @@ joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
   (void)last;
   (void)result;
   (void)init;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -897,8 +898,8 @@ inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
       g, typename sycl::detail::GroupOpTag<T>::type(), x, binary_op);
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -919,8 +920,8 @@ inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
   (void)g;
   (void)x;
   (void)binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -946,8 +947,8 @@ inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
   return inclusive_scan_over_group(g, y, binary_op);
 #else
   (void)g;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -1015,8 +1016,8 @@ joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
   (void)g;
   (void)last;
   (void)result;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Group algorithms are not supported on host.");
 #endif
 }
 
diff --git a/sycl/include/sycl/group_barrier.hpp b/sycl/include/sycl/group_barrier.hpp
index c91dddd427aa6..43414f3125599 100644
--- a/sycl/include/sycl/group_barrier.hpp
+++ b/sycl/include/sycl/group_barrier.hpp
@@ -14,6 +14,7 @@
 #include <CL/__spirv/spirv_vars.hpp>
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/group.hpp>
 #include <sycl/sub_group.hpp>
 
@@ -31,8 +32,8 @@ group_barrier(Group G, memory_scope FenceScope = Group::fence_scope) {
 #else
   (void)G;
   (void)FenceScope;
-  throw sycl::runtime_error("Barriers are not supported on host device",
-                            PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Barriers are not supported on host device");
 #endif
 }
 
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 91f3a9aa336fd..5e42f3faad078 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -197,7 +197,7 @@ checkValueRangeImpl(ValT V) {
   static constexpr size_t Limit =
       static_cast<size_t>((std::numeric_limits<int>::max)());
   if (V > Limit)
-    throw runtime_error(NotIntMsg<T>::Msg, PI_ERROR_INVALID_VALUE);
+    throw sycl::exception(make_error_code(errc::nd_range), NotIntMsg<T>::Msg);
 }
 #endif
 
@@ -372,10 +372,10 @@ class __SYCL_EXPORT handler {
 
   void throwIfActionIsCreated() {
     if (detail::CG::None != getType())
-      throw sycl::runtime_error("Attempt to set multiple actions for the "
-                                "command group. Command group must consist of "
-                                "a single kernel or explicit memory operation.",
-                                PI_ERROR_INVALID_OPERATION);
+      throw sycl::exception(make_error_code(errc::runtime),
+                            "Attempt to set multiple actions for the "
+                            "command group. Command group must consist of "
+                            "a single kernel or explicit memory operation.");
   }
 
   constexpr static int AccessTargetMask = 0x7ff;
diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
index add48de0fa8b1..27e7144851fa9 100644
--- a/sycl/include/sycl/queue.hpp
+++ b/sycl/include/sycl/queue.hpp
@@ -18,6 +18,7 @@
 #include <sycl/device.hpp>
 #include <sycl/device_selector.hpp>
 #include <sycl/event.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/exception_list.hpp>
 #include <sycl/ext/oneapi/device_global/device_global.hpp>
 #include <sycl/ext/oneapi/weak_object_base.hpp>
@@ -2345,9 +2346,9 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
       // which won't be properly resolved in separate compile use-case
 #ifndef NDEBUG
       if (AH->Flag == __SYCL_ASSERT_START)
-        throw sycl::runtime_error(
-            "Internal Error. Invalid value in assert description.",
-            PI_ERROR_INVALID_VALUE);
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "Internal Error. Invalid value in assert description.");
 #endif
 
       if (AH->Flag) {
diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp
index ed82fb44c6f0a..31dcf80f7ec2b 100644
--- a/sycl/include/sycl/reduction.hpp
+++ b/sycl/include/sycl/reduction.hpp
@@ -12,6 +12,7 @@
 #include <sycl/atomic.hpp>
 #include <sycl/atomic_ref.hpp>
 #include <sycl/detail/tuple.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/accessor_property_list.hpp>
 #include <sycl/group_algorithm.hpp>
 #include <sycl/handler.hpp>
@@ -1109,15 +1110,13 @@ class reduction_impl
       : algo(BOp, InitializeToIdentity, Var) {
     if constexpr (!is_usm)
       if (Var.size() != 1)
-        throw sycl::runtime_error(errc::invalid,
-                                  "Reduction variable must be a scalar.",
-                                  PI_ERROR_INVALID_VALUE);
+        throw sycl::exception(make_error_code(errc::invalid),
+                              "Reduction variable must be a scalar.");
     if constexpr (!is_known_identity)
       if (InitializeToIdentity)
-        throw sycl::runtime_error(errc::invalid,
-                                  "initialize_to_identity property cannot be "
-                                  "used with identityless reductions.",
-                                  PI_ERROR_INVALID_VALUE);
+        throw sycl::exception(make_error_code(errc::invalid),
+                              "initialize_to_identity property cannot be "
+                              "used with identityless reductions.");
   }
 
   /// Constructs reduction_impl with an explicit identity value. This is only
@@ -1129,9 +1128,8 @@ class reduction_impl
       : algo(Identity, BOp, InitializeToIdentity, Var) {
     if constexpr (!is_usm)
       if (Var.size() != 1)
-        throw sycl::runtime_error(errc::invalid,
-                                  "Reduction variable must be a scalar.",
-                                  PI_ERROR_INVALID_VALUE);
+        throw sycl::exception(make_error_code(errc::invalid),
+                              "Reduction variable must be a scalar.");
   }
 };
 
@@ -1680,11 +1678,11 @@ struct NDRangeReduction<
     // for the device.
     size_t MaxWGSize = reduGetMaxWGSize(Queue, OneElemSize);
     if (NDRange.get_local_range().size() > MaxWGSize)
-      throw sycl::runtime_error("The implementation handling parallel_for with"
-                                " reduction requires work group size not bigger"
-                                " than " +
-                                    std::to_string(MaxWGSize),
-                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+      throw sycl::exception(make_error_code(errc::nd_range),
+                            "The implementation handling parallel_for with"
+                            " reduction requires work group size not bigger"
+                            " than " +
+                                std::to_string(MaxWGSize));
 
     size_t NElements = Reduction::num_elements;
     size_t NWorkGroups = NDRange.get_group_range().size();
@@ -1725,13 +1723,13 @@ struct NDRangeReduction<
     // TODO: Create a special slow/sequential version of the kernel that would
     // handle the reduction instead of reporting an assert below.
     if (MaxWGSize <= 1)
-      throw sycl::runtime_error("The implementation handling parallel_for with "
-                                "reduction requires the maximal work group "
-                                "size to be greater than 1 to converge. "
-                                "The maximal work group size depends on the "
-                                "device and the size of the objects passed to "
-                                "the reduction.",
-                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+      throw sycl::exception(make_error_code(errc::nd_range),
+                            "The implementation handling parallel_for with "
+                            "reduction requires the maximal work group "
+                            "size to be greater than 1 to converge. "
+                            "The maximal work group size depends on the "
+                            "device and the size of the objects passed to "
+                            "the reduction.");
     size_t NWorkItems = NDRange.get_group_range().size();
     while (NWorkItems > 1) {
       reduction::withAuxHandler(CGH, [&](handler &AuxHandler) {
@@ -1808,11 +1806,11 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
     // compiled for the device.
     size_t MaxWGSize = reduGetMaxWGSize(Queue, OneElemSize);
     if (NDRange.get_local_range().size() > MaxWGSize)
-      throw sycl::runtime_error("The implementation handling parallel_for with"
-                                " reduction requires work group size not bigger"
-                                " than " +
-                                    std::to_string(MaxWGSize),
-                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+      throw sycl::exception(make_error_code(errc::nd_range),
+                            "The implementation handling parallel_for with"
+                            " reduction requires work group size not bigger"
+                            " than " +
+                                std::to_string(MaxWGSize));
 
     size_t NWorkGroups = NDRange.get_group_range().size();
 
@@ -1904,13 +1902,13 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
     // TODO: Create a special slow/sequential version of the kernel that would
     // handle the reduction instead of reporting an assert below.
     if (MaxWGSize <= 1)
-      throw sycl::runtime_error("The implementation handling parallel_for with "
-                                "reduction requires the maximal work group "
-                                "size to be greater than 1 to converge. "
-                                "The maximal work group size depends on the "
-                                "device and the size of the objects passed to "
-                                "the reduction.",
-                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+      throw sycl::exception(make_error_code(errc::nd_range),
+                            "The implementation handling parallel_for with "
+                            "reduction requires the maximal work group "
+                            "size to be greater than 1 to converge. "
+                            "The maximal work group size depends on the "
+                            "device and the size of the objects passed to "
+                            "the reduction.");
     size_t NWorkItems = NDRange.get_group_range().size();
     while (NWorkItems > 1) {
       size_t NWorkGroups;
@@ -2586,11 +2584,11 @@ template <> struct NDRangeReduction<reduction::strategy::multi> {
     // for the device.
     size_t MaxWGSize = reduGetMaxWGSize(Queue, LocalMemPerWorkItem);
     if (NDRange.get_local_range().size() > MaxWGSize)
-      throw sycl::runtime_error("The implementation handling parallel_for with"
-                                " reduction requires work group size not bigger"
-                                " than " +
-                                    std::to_string(MaxWGSize),
-                                PI_ERROR_INVALID_WORK_GROUP_SIZE);
+      throw sycl::exception(make_error_code(errc::nd_range),
+                            "The implementation handling parallel_for with"
+                            " reduction requires work group size not bigger"
+                            " than " +
+                                std::to_string(MaxWGSize));
 
     reduCGFuncMulti<KernelName>(CGH, KernelFunc, NDRange, Properties, ReduTuple,
                                 ReduIndices);
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
index daf4294234dec..6c0cd517e3aca 100644
--- a/sycl/include/sycl/sub_group.hpp
+++ b/sycl/include/sycl/sub_group.hpp
@@ -8,6 +8,7 @@
 
 #pragma once
 
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/sub_group.hpp>
 #include <sycl/group.hpp>
 
@@ -22,8 +23,8 @@ inline sub_group this_sub_group() {
 #ifdef __SYCL_DEVICE_ONLY__
   return sub_group();
 #else
-  throw runtime_error("Sub-groups are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
 #endif
 }
 } // namespace ext::oneapi::experimental
diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp
index 8a334e2e56b4c..dc26dc249cb81 100644
--- a/sycl/include/sycl/types.hpp
+++ b/sycl/include/sycl/types.hpp
@@ -51,6 +51,7 @@
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/type_traits.hpp>
+#include <sycl/exception.hpp>
 #include <sycl/half_type.hpp>
 #include <sycl/marray.hpp>
 #include <sycl/multi_ptr.hpp>
@@ -310,12 +311,13 @@ std::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",
-                          PI_ERROR_UNKNOWN);
+      throw sycl::exception(make_error_code(errc::runtime),
+                            "Unable to set rounding mode to FE_TONEAREST");
     R Result = std::rint(Value);
     Err = std::fesetround(OldRoundingDirection);
     if (Err)
-      throw runtime_error("Unable to restore rounding mode.", PI_ERROR_UNKNOWN);
+      throw sycl::exception(make_error_code(errc::runtime),
+                            "Unable to restore rounding mode.");
     return Result;
   }
     // Round toward zero.
diff --git a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp
index 9550eb84310b3..718f92f9270e9 100644
--- a/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp
+++ b/sycl/test-e2e/Basic/range_offset_fit_in_int.cpp
@@ -8,7 +8,7 @@
 
 namespace S = sycl;
 
-void checkRangeException(S::runtime_error &E) {
+void checkRangeException(S::exception &E) {
   constexpr char Msg[] = "Provided range is out of integer limits. "
                          "Pass `-fno-sycl-id-queries-fit-in-int' to "
                          "disable range check.";
@@ -18,7 +18,7 @@ void checkRangeException(S::runtime_error &E) {
   assert(std::string(E.what()).find(Msg) == 0 && "Unexpected message");
 }
 
-void checkOffsetException(S::runtime_error &E) {
+void checkOffsetException(S::exception &E) {
   constexpr char Msg[] = "Provided offset is out of integer limits. "
                          "Pass `-fno-sycl-id-queries-fit-in-int' to "
                          "disable offset check.";
@@ -73,7 +73,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -101,7 +101,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -117,7 +117,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -134,7 +134,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -150,7 +150,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkOffsetException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -178,7 +178,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -194,7 +194,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -210,7 +210,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkOffsetException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -238,7 +238,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
@@ -254,7 +254,7 @@ void test() {
     });
 
     assert(false && "Exception expected");
-  } catch (S::runtime_error &E) {
+  } catch (S::exception &E) {
     checkRangeException(E);
   } catch (...) {
     assert(false && "Unexpected exception catched");
diff --git a/sycl/test/basic_tests/group.cpp b/sycl/test/basic_tests/group.cpp
index 06872cbfc2d2f..c72eef464b625 100644
--- a/sycl/test/basic_tests/group.cpp
+++ b/sycl/test/basic_tests/group.cpp
@@ -35,13 +35,13 @@ int main() {
     one_dim.get_local_id();
     assert(one_dim.get_local_id(0) == one_dim.get_local_id()[0]);
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 
   try {
     one_dim.get_local_linear_id();
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 
   // two dimension group
@@ -68,13 +68,13 @@ int main() {
     assert(two_dim.get_local_id(0) == two_dim.get_local_id()[0]);
     assert(two_dim.get_local_id(1) == two_dim.get_local_id()[1]);
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 
   try {
     two_dim.get_local_linear_id();
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 
   // three dimension group
@@ -108,12 +108,12 @@ int main() {
     assert(three_dim.get_local_id(1) == three_dim.get_local_id()[1]);
     assert(three_dim.get_local_id(2) == three_dim.get_local_id()[2]);
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 
   try {
     three_dim.get_local_linear_id();
     assert(0); // get_local_id() is not implemented on host device
-  } catch (sycl::runtime_error) {
+  } catch (sycl::exception) {
   }
 }