diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
index 651dd5f578234..ef4ae2995e29d 100644
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -10746,11 +10746,6 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
   if (TCArgs.hasFlag(options::OPT_fno_sycl_esimd_force_stateless_mem,
                      options::OPT_fsycl_esimd_force_stateless_mem, false))
     addArgs(PostLinkArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"});
-
-  bool IsUsingLTO = TC.getDriver().isUsingLTO(/*IsDeviceOffloadAction=*/true);
-  auto LTOMode = TC.getDriver().getLTOMode(/*IsDeviceOffloadAction=*/true);
-  if (!IsUsingLTO || LTOMode != LTOK_Thin)
-    addArgs(PostLinkArgs, TCArgs, {"-properties"});
 }
 
 // Add any sycl-post-link options that rely on a specific Triple in addition
@@ -10765,11 +10760,18 @@ static void getTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
                                            llvm::Triple Triple,
                                            bool SpecConstsSupported,
                                            types::ID OutputType) {
+
+  bool IsUsingLTO = TC.getDriver().isUsingLTO(/*IsDeviceOffloadAction=*/true);
+  auto LTOMode = TC.getDriver().getLTOMode(/*IsDeviceOffloadAction=*/true);
   if (OutputType == types::TY_LLVM_BC) {
     // single file output requested - this means only perform necessary IR
     // transformations (like specialization constant intrinsic lowering) and
     // output LLVMIR
     addArgs(PostLinkArgs, TCArgs, {"-ir-output-only"});
+  } else if (!IsUsingLTO || LTOMode != LTOK_Thin) {
+    // Only create a properties file if we are not
+    // only outputting IR.
+    addArgs(PostLinkArgs, TCArgs, {"-properties"});
   }
   if (SpecConstsSupported)
     addArgs(PostLinkArgs, TCArgs, {"-spec-const=native"});
diff --git a/clang/test/Driver/sycl-offload-new-driver.c b/clang/test/Driver/sycl-offload-new-driver.c
index ded7e89252e7c..be71376bfacda 100644
--- a/clang/test/Driver/sycl-offload-new-driver.c
+++ b/clang/test/Driver/sycl-offload-new-driver.c
@@ -60,7 +60,7 @@
 // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
 // RUN:          -Xdevice-post-link -post-link-opt -### %s 2>&1 \
 // RUN:   | FileCheck -check-prefix WRAPPER_OPTIONS_POSTLINK %s
-// WRAPPER_OPTIONS_POSTLINK: clang-linker-wrapper{{.*}} "--sycl-post-link-options=-O2 -device-globals -properties -post-link-opt"
+// WRAPPER_OPTIONS_POSTLINK: clang-linker-wrapper{{.*}} "--sycl-post-link-options=-O2 -device-globals -post-link-opt"
 
 // -fsycl-device-only behavior
 // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \
diff --git a/clang/test/Driver/sycl-post-link-options-win.cpp b/clang/test/Driver/sycl-post-link-options-win.cpp
index 110dcfb498163..7c641cea6783d 100644
--- a/clang/test/Driver/sycl-post-link-options-win.cpp
+++ b/clang/test/Driver/sycl-post-link-options-win.cpp
@@ -20,4 +20,4 @@
 // RUN:   -sycl-device-libraries=%t.devicelib.o \
 // RUN:   --sycl-post-link-options="-O2 -device-globals -O0" \
 // RUN:   --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck --check-prefix OPTIONS_POSTLINK_JIT_NEW %s
-// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -O0
+// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -properties -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -O0
diff --git a/clang/test/Driver/sycl-post-link-options.cpp b/clang/test/Driver/sycl-post-link-options.cpp
index 96134faba4ae2..1f1f4d5a7c07f 100644
--- a/clang/test/Driver/sycl-post-link-options.cpp
+++ b/clang/test/Driver/sycl-post-link-options.cpp
@@ -18,6 +18,6 @@
 //
 // RUN: clang-linker-wrapper --dry-run --host-triple=x86_64-unknown-linux-gnu \
 // RUN:   -sycl-device-libraries=%t.devicelib.o \
-// RUN:   --sycl-post-link-options="-O2 -device-globals -properties -O0" \
+// RUN:   --sycl-post-link-options="-O2 -device-globals -O0" \
 // RUN:   --linker-path=/usr/bin/ld %t.o -o a.out 2>&1 | FileCheck --check-prefix OPTIONS_POSTLINK_JIT_NEW %s
-// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -properties -O0
+// OPTIONS_POSTLINK_JIT_NEW: sycl-post-link{{.*}} -spec-const=native -properties -split=auto -emit-only-kernels-as-entry-points -emit-param-info -symbols -emit-exported-symbols -emit-imported-symbols -split-esimd -lower-esimd -O2 -device-globals -O0
diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
index c128179cdb323..c0def22c4147f 100644
--- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
+++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp
@@ -545,6 +545,9 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args,
   else
     PostLinkArgs.push_back("-spec-const=emulation");
 
+  if (Triple.isSPIROrSPIRV() || SYCLNativeCPU)
+    PostLinkArgs.push_back("-properties");
+
   // See if device code splitting is already requested. If not requested, then
   // set -split=auto for non-FPGA targets.
   bool NoSplit = true;
diff --git a/sycl/include/sycl/ext/oneapi/backend/hip.hpp b/sycl/include/sycl/ext/oneapi/backend/hip.hpp
index 904d4c5582e32..0f59dd2f4116a 100644
--- a/sycl/include/sycl/ext/oneapi/backend/hip.hpp
+++ b/sycl/include/sycl/ext/oneapi/backend/hip.hpp
@@ -17,10 +17,8 @@ inline namespace _V1 {
 template <>
 inline backend_return_t<backend::ext_oneapi_hip, device>
 get_native<backend::ext_oneapi_hip, device>(const device &Obj) {
-  // TODO swap with SYCL 2020 exception when in ABI-break window
   if (Obj.get_backend() != backend::ext_oneapi_hip) {
-    throw sycl::runtime_error(errc::backend_mismatch, "Backends mismatch",
-                              PI_ERROR_INVALID_OPERATION);
+    throw exception(errc::backend_mismatch, "Backends mismatch");
   }
   // HIP 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/ext/oneapi/bf16_storage_builtins.hpp b/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
index 0cb67f9a7d377..b0935475d63a2 100644
--- a/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
+++ b/sycl/include/sycl/ext/oneapi/bf16_storage_builtins.hpp
@@ -49,8 +49,8 @@ std::enable_if_t<detail::is_bf16_storage_type<T>::value, T> fabs(T x) {
   return __clc_fabs(x);
 #else
   (void)x;
-  throw runtime_error("bf16 is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "bf16 is not supported on host.");
 #endif
 }
 template <typename T>
@@ -60,8 +60,8 @@ std::enable_if_t<detail::is_bf16_storage_type<T>::value, T> fmin(T x, T y) {
 #else
   (void)x;
   (void)y;
-  throw runtime_error("bf16 is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "bf16 is not supported on host.");
 #endif
 }
 template <typename T>
@@ -71,8 +71,8 @@ std::enable_if_t<detail::is_bf16_storage_type<T>::value, T> fmax(T x, T y) {
 #else
   (void)x;
   (void)y;
-  throw runtime_error("bf16 is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "bf16 is not supported on host.");
 #endif
 }
 template <typename T>
@@ -83,8 +83,8 @@ std::enable_if_t<detail::is_bf16_storage_type<T>::value, T> fma(T x, T y, T z) {
   (void)x;
   (void)y;
   (void)z;
-  throw runtime_error("bf16 is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "bf16 is not supported on host.");
 #endif
 }
 
diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp
index bbee229aaf4ea..9fb099e7eb8e4 100644
--- a/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp
+++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-intel.hpp
@@ -15,8 +15,7 @@
 #include <sycl/access/access.hpp>             // for address_space, decorated
 #include <sycl/builtins.hpp>                  // for fabs
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
-#include <sycl/detail/pi.h>                   // for PI_ERROR_INVALID_DEVICE
-#include <sycl/exception.hpp>                 // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/bfloat16.hpp>       // for bfloat16
 #include <sycl/ext/oneapi/experimental/annotated_ptr/annotated_ptr.hpp> // for annotated_ptr
 #include <sycl/group.hpp>     // for group
@@ -122,8 +121,8 @@ class wi_element {
     const size_t col = coord[1];
     return std::make_tuple(row, col);
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -137,8 +136,8 @@ class wi_element {
                                                                      idx);
     return elem;
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -151,8 +150,8 @@ class wi_element {
                                         spv_scope_traits<Group>::value>(
                M.spvm, idx) != static_cast<storage_element_type>(0);
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -163,8 +162,8 @@ class wi_element {
     return *this;
 #else
     (void)rhs;
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -182,8 +181,8 @@ class wi_element {
     return *this;
 #else
     (void)rhs;
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -206,8 +205,8 @@ class wi_element {
 #define OP(op)                                                                 \
   template <typename T2> wi_element &operator op##=(const T2 & rhs) {          \
     (void)rhs;                                                                 \
-    throw runtime_error("joint matrix is not supported on host device.",       \
-                        PI_ERROR_INVALID_DEVICE);                              \
+    throw exception(make_error_code(errc::runtime),                            \
+                    "joint matrix is not supported on host.");                 \
   }
 #endif // __SYCL_DEVICE_ONLY__
   OP(+)
@@ -242,8 +241,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
     const uint32_t col = coord[1];
     return std::make_tuple(row, col);
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -255,8 +254,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
         spv_matrix_layout_traits<Layout>::value,
         spv_scope_traits<Group>::value>(M.spvm, idx);
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -270,8 +269,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
                    spv_scope_traits<Group>::value>(M.spvm, idx))) >=
            std::numeric_limits<float>::epsilon();
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -281,8 +280,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
     return *this;
 #else
     (void)rhs;
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -301,8 +300,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
     return *this;
 #else
     (void)rhs;
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif // __SYCL_DEVICE_ONLY__
   }
 
@@ -323,8 +322,8 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
 #define OP(opassign, op)                                                       \
   wi_element &operator opassign(const sycl::ext::oneapi::bfloat16 & rhs) {     \
     (void)rhs;                                                                 \
-    throw runtime_error("joint matrix is not supported on host device.",       \
-                        PI_ERROR_INVALID_DEVICE);                              \
+    throw exception(make_error_code(errc::runtime),                            \
+                    "joint matrix is not supported on host.");                 \
   }
 #endif // __SYCL_DEVICE_ONLY__
   OP(+=, +)
@@ -398,15 +397,15 @@ class wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use, Layout,
       const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use,     \
                        Layout, Group> &,                                       \
       const sycl::ext::oneapi::bfloat16 &) {                                   \
-    throw runtime_error("joint matrix is not supported on host device.",       \
-                        PI_ERROR_INVALID_DEVICE);                              \
+    throw exception(make_error_code(errc::runtime),                            \
+                    "joint matrix is not supported on host.");                 \
   }                                                                            \
   friend type operator op(                                                     \
       const sycl::ext::oneapi::bfloat16 &,                                     \
       const wi_element<sycl::ext::oneapi::bfloat16, NumRows, NumCols, Use,     \
                        Layout, Group> &) {                                     \
-    throw runtime_error("joint matrix is not supported on host device.",       \
-                        PI_ERROR_INVALID_DEVICE);                              \
+    throw exception(make_error_code(errc::runtime),                            \
+                    "joint matrix is not supported on host.");                 \
   }
   OP(sycl::ext::oneapi::bfloat16, +)
   OP(sycl::ext::oneapi::bfloat16, -)
@@ -451,8 +450,8 @@ class wi_data {
 #if __SYCL_DEVICE_ONLY__
     return __spirv_JointMatrixWorkItemLengthINTEL(jm.spvm);
 #else
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif
   };
 
@@ -496,10 +495,10 @@ joint_matrix_store(Group,
   std::ignore = src;
   std::ignore = dst;
   std::ignore = stride;
-  throw runtime_error(
+  throw exception(
+      make_error_code(errc::runtime),
       "This version of the matrix extension is only currently supported on "
-      "intel devices",
-      PI_ERROR_INVALID_DEVICE);
+      "intel devices");
 #else
   // intel's impl
   using DecorT = typename sycl::detail::DecoratedType<T, Space>::type;
@@ -518,8 +517,8 @@ joint_matrix_store(Group,
   std::ignore = src;
   std::ignore = dst;
   std::ignore = stride;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -542,10 +541,10 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
   std::ignore = src;
   std::ignore = dst;
   std::ignore = stride;
-  throw runtime_error(
+  throw exception(
+      make_error_code(errc::runtime),
       "This version of the matrix extension is only currently supported on "
-      "intel devices",
-      PI_ERROR_INVALID_DEVICE);
+      "intel devices");
 #else
   // intel's impl
   T *Ptr = dst.get();
@@ -563,8 +562,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
   std::ignore = src;
   std::ignore = dst;
   std::ignore = stride;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -599,8 +598,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_apply(
   std::ignore = sg;
   std::ignore = jm;
   std::ignore = lambda;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif
 }
 
@@ -630,8 +629,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_fill_checked(
   std::ignore = Width;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -668,8 +667,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load_checked(
   std::ignore = Layout;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -705,8 +704,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load_checked(
   std::ignore = Width;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -740,8 +739,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store_checked(
   std::ignore = Layout;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -773,8 +772,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store_checked(
   std::ignore = Width;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -808,8 +807,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load_checked(
   std::ignore = Layout;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -841,8 +840,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load_checked(
   std::ignore = Width;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -873,8 +872,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store_checked(
   std::ignore = Layout;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -902,8 +901,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store_checked(
   std::ignore = Width;
   std::ignore = CoordX;
   std::ignore = CoordY;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 // End out-of-bounds API
diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
index c8e473217ef48..cba25494ca65d 100644
--- a/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
+++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-unified.hpp
@@ -20,8 +20,7 @@
 
 #include <sycl/access/access.hpp>             // for address_space
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_...
-#include <sycl/detail/pi.h>                   // for PI_ERROR_INVAL...
-#include <sycl/exception.hpp>                 // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/matrix/matrix-unified-utils.hpp> // for layout, use, tf32, convertMatrixUseEnumToString
 #include <sycl/ext/oneapi/matrix/query-types.hpp> // for convertTypeToMatrixTypeString
 #include <sycl/marray.hpp>                        // for marray
@@ -69,8 +68,8 @@ struct joint_matrix {
 #endif // defined(__SYCL_DEVICE_ONLY__)
   joint_matrix() {
 #ifndef __SYCL_DEVICE_ONLY__
-    throw runtime_error("joint matrix is not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "joint matrix is not supported on host.");
 #endif
   }
 #ifdef __SYCL_DEVICE_ONLY__
@@ -107,8 +106,8 @@ joint_matrix_apply(Group sg, joint_matrix<Group, T, Use, M, N, Layout> &jm,
   std::ignore = sg;
   std::ignore = jm;
   std::ignore = lambda;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif
   return;
 }
@@ -143,8 +142,8 @@ joint_matrix_apply(Group sg, joint_matrix<Group, T, Use, M, N, Layout> &jmsrc,
   std::ignore = jmsrc;
   std::ignore = jmdest;
   std::ignore = lambda;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif
   return;
 }
@@ -171,8 +170,8 @@ joint_matrix_fill(Group,
 #else
   std::ignore = res;
   std::ignore = v;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -214,8 +213,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
   std::ignore = src;
   std::ignore = stride;
   std::ignore = Layout;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -259,8 +258,8 @@ joint_matrix_load(Group sg,
   std::ignore = res;
   std::ignore = src;
   std::ignore = stride;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -277,11 +276,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
 #if defined(__SYCL_DEVICE_ONLY__)
 #if defined(__NVPTX__)
   std::ignore = sg;
-  throw runtime_error("Use joint_matrix_load on multi_ptr on Nvidia device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_load on multi_ptr on Nvidia device.");
 #elif defined(__HIP_PLATFORM_AMD_MFMA__)
-  throw runtime_error("Use joint_matrix_load on multi_ptr on AMD device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_load on multi_ptr on AMD device.");
 #else
   std::ignore = sg;
   T *Ptr = src.get();
@@ -297,8 +296,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
   std::ignore = src;
   std::ignore = stride;
   std::ignore = Layout;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -316,11 +315,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
 #if defined(__SYCL_DEVICE_ONLY__)
 #if defined(__NVPTX__)
   std::ignore = sg;
-  throw runtime_error("Use joint_matrix_load on multi_ptr on Nvidia device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_load on multi_ptr on Nvidia device.");
 #elif defined(__HIP_PLATFORM_AMD_MFMA__)
-  throw runtime_error("Use joint_matrix_load on multi_ptr on AMD device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_load on multi_ptr on AMD device.");
 #else
   std::ignore = sg;
   T *Ptr = src.get();
@@ -336,8 +335,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_load(
   std::ignore = res;
   std::ignore = src;
   std::ignore = stride;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -379,8 +378,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
   std::ignore = dst;
   std::ignore = stride;
   std::ignore = Layout;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -396,11 +395,11 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
 #if defined(__SYCL_DEVICE_ONLY__)
 #if defined(__NVPTX__)
   std::ignore = sg;
-  throw runtime_error("Use joint_matrix_store on multi_ptr on Nvidia device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_store on multi_ptr on Nvidia device.");
 #elif defined(__HIP_PLATFORM_AMD_MFMA__)
-  throw runtime_error("Use joint_matrix_store on multi_ptr on AMD device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Use joint_matrix_store on multi_ptr on AMD device.");
 #else
   std::ignore = sg;
   T *Ptr = dst.get();
@@ -416,8 +415,8 @@ inline __SYCL_ALWAYS_INLINE void joint_matrix_store(
   std::ignore = dst;
   std::ignore = stride;
   std::ignore = Layout;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -483,8 +482,8 @@ joint_matrix_mad(
   std::ignore = B;
   std::ignore = C;
   std::ignore = D;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -511,8 +510,8 @@ void joint_matrix_copy(
   std::ignore = sg;
   std::ignore = dst;
   std::ignore = src;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
@@ -546,14 +545,13 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride,
 #if defined(__NVPTX__)
   std::ignore = sg;
   std::ignore = properties;
-  throw runtime_error(
-      "joint_matrix_prefetch is not supported on Nvidia device.",
-      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint_matrix_prefetch is not supported on Nvidia device.");
 #elif defined(__HIP_PLATFORM_AMD_MFMA__)
   std::ignore = sg;
   std::ignore = properties;
-  throw runtime_error("joint_matrix_prefetch is not supported on AMD device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint_matrix_prefetch is not supported on AMD device.");
 #else
   std::ignore = sg;
   auto prop = properties.template get_property<prefetch_hint_key>();
@@ -567,8 +565,8 @@ joint_matrix_prefetch(Group sg, T *Ptr, size_t stride,
   std::ignore = stride;
   std::ignore = Layout;
   std::ignore = properties;
-  throw runtime_error("joint matrix is not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "joint matrix is not supported on host.");
 #endif // defined(__SYCL_DEVICE_ONLY__)
 }
 
diff --git a/sycl/include/sycl/vector_preview.hpp b/sycl/include/sycl/vector_preview.hpp
index d0549dbefd817..e374e0309445c 100644
--- a/sycl/include/sycl/vector_preview.hpp
+++ b/sycl/include/sycl/vector_preview.hpp
@@ -489,9 +489,9 @@ class vec : public detail::vec_arith<DataT, NumElements> {
           !std::is_same_v<convertT, bool>;
 
       if constexpr (canUseNativeVectorConvert) {
-        Result.m_Data = sycl::bit_cast<decltype(Result.m_Data)>(
-            detail::convertImpl<T, R, roundingMode, NumElements, OpenCLVecT,
-                                OpenCLVecR>(NativeVector));
+        auto val = detail::convertImpl<T, R, roundingMode, NumElements, OpenCLVecT,
+                                OpenCLVecR>(NativeVector);
+        Result.m_Data = sycl::bit_cast<decltype(Result.m_Data)>(val);
       } else
 #endif // __SYCL_DEVICE_ONLY__
       {
diff --git a/sycl/plugins/unified_runtime/CMakeLists.txt b/sycl/plugins/unified_runtime/CMakeLists.txt
index 359989be5fd26..2cf69449e95ec 100644
--- a/sycl/plugins/unified_runtime/CMakeLists.txt
+++ b/sycl/plugins/unified_runtime/CMakeLists.txt
@@ -105,7 +105,7 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   # Date:   Mon May 20 15:50:02 2024 +0100
   #     Merge pull request #954 from jchlanda/jakub/rqwgs_hip
   #     [HIP] Handle required wg size attribute in HIP
-  set(UNIFIED_RUNTIME_TAG b167fa1b331b01dc0d904e37e98dc240848c244f)
+  set(UNIFIED_RUNTIME_TAG e1af8f87e11927c9e8fcc3bbbf0d5097e864359c)
 
   fetch_adapter_source(level_zero
     ${UNIFIED_RUNTIME_REPO}
diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp
index cb142906a5970..ad00aafe22ac2 100644
--- a/sycl/source/backend.cpp
+++ b/sycl/source/backend.cpp
@@ -61,8 +61,8 @@ backend convertBackend(pi_platform_backend PiBackend) {
   case PI_EXT_PLATFORM_BACKEND_NATIVE_CPU:
     return backend::ext_oneapi_native_cpu;
   }
-  throw sycl::runtime_error{"convertBackend: Unsupported backend",
-                            PI_ERROR_INVALID_OPERATION};
+  throw exception(make_error_code(errc::runtime),
+                  "convertBackend: Unsupported backend");
 }
 
 platform make_platform(pi_native_handle NativeHandle, backend Backend) {
diff --git a/sycl/source/detail/config.hpp b/sycl/source/detail/config.hpp
index e94be9bd27166..4659805ca76fa 100644
--- a/sycl/source/detail/config.hpp
+++ b/sycl/source/detail/config.hpp
@@ -381,7 +381,7 @@ template <> class SYCLConfig<SYCL_CACHE_PERSISTENT> {
       std::string Msg =
           std::string{"Invalid value for bool configuration variable "} +
           getName() + std::string{": "} + ValStr;
-      throw runtime_error(Msg, PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime), Msg);
     }
     return ValStr[0] == '1';
   }
@@ -603,7 +603,7 @@ template <> class SYCLConfig<SYCL_CACHE_IN_MEM> {
       std::string Msg =
           std::string{"Invalid value for bool configuration variable "} +
           getName() + std::string{": "} + ValStr;
-      throw runtime_error(Msg, PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime), Msg);
     }
     return ValStr[0] == '1';
   }
diff --git a/sycl/source/detail/error_handling/error_handling.cpp b/sycl/source/detail/error_handling/error_handling.cpp
index eb95d979560db..6657620f5ed61 100644
--- a/sycl/source/detail/error_handling/error_handling.cpp
+++ b/sycl/source/detail/error_handling/error_handling.cpp
@@ -301,10 +301,8 @@ void handleInvalidWorkGroupSize(const device_impl &DeviceImpl, pi_kernel Kernel,
   // consistent with the required number of sub-groups for kernel in the
   // program source.
 
-  // Fallback
-  constexpr pi_result Error = PI_ERROR_INVALID_WORK_GROUP_SIZE;
-  throw runtime_error(
-      "PI backend failed. PI backend returns: " + codeToString(Error), Error);
+  throw exception(make_error_code(errc::nd_range),
+                  "internal error: expected HasLocalSize");
 }
 
 void handleInvalidWorkItemSize(const device_impl &DeviceImpl,
@@ -348,9 +346,7 @@ void handleInvalidValue(const device_impl &DeviceImpl,
   }
 
   // fallback
-  constexpr pi_result Error = PI_ERROR_INVALID_VALUE;
-  throw runtime_error(
-      "Native API failed. Native API returns: " + codeToString(Error), Error);
+  throw exception(make_error_code(errc::nd_range), "unknown internal error");
 }
 
 void handleErrorOrWarning(pi_result Error, const device_impl &DeviceImpl,
@@ -424,8 +420,8 @@ void handleErrorOrWarning(pi_result Error, const device_impl &DeviceImpl,
     // TODO: Handle other error codes
 
   default:
-    throw runtime_error(
-        "Native API failed. Native API returns: " + codeToString(Error), Error);
+    throw detail::set_pi_error(
+        exception(make_error_code(errc::runtime), "PI error"), Error);
   }
 }
 
diff --git a/sycl/source/detail/filter_selector_impl.cpp b/sycl/source/detail/filter_selector_impl.cpp
index 0043622d62483..69b53c05d21b4 100644
--- a/sycl/source/detail/filter_selector_impl.cpp
+++ b/sycl/source/detail/filter_selector_impl.cpp
@@ -56,7 +56,7 @@ filter create_filter(const std::string &Input) {
   // There should only be up to 3 tokens.
   // BE:Device Type:Device Num
   if (Tokens.size() > 3)
-    throw sycl::runtime_error(Error, PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid), Error);
 
   for (const std::string &Token : Tokens) {
     if (Token == "cpu" && !Result.DeviceType) {
@@ -77,10 +77,10 @@ filter create_filter(const std::string &Input) {
       try {
         Result.DeviceNum = std::stoi(Token);
       } catch (std::logic_error &) {
-        throw sycl::runtime_error(Error, PI_ERROR_INVALID_VALUE);
+        throw exception(make_error_code(errc::invalid), Error);
       }
     } else {
-      throw sycl::runtime_error(Error, PI_ERROR_INVALID_VALUE);
+      throw exception(make_error_code(errc::invalid), Error);
     }
   }
 
@@ -141,9 +141,9 @@ int filter_selector_impl::operator()(const device &Dev) const {
 
   mNumDevicesSeen++;
   if ((mNumDevicesSeen == mNumTotalDevices) && !mMatchFound) {
-    throw sycl::runtime_error(
-        "Could not find a device that matches the specified filter(s)!",
-        PI_ERROR_DEVICE_NOT_FOUND);
+    throw exception(
+        make_error_code(errc::runtime),
+        "Could not find a device that matches the specified filter(s)!");
   }
 
   return Score;
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
index 35a5b690ccb8a..a532e5504a09f 100644
--- a/sycl/source/detail/kernel_bundle_impl.hpp
+++ b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -152,9 +152,9 @@ class kernel_bundle_impl {
         break;
       case bundle_state::input:
       case bundle_state::ext_oneapi_source:
-        throw sycl::runtime_error("Internal error. The target state should not "
-                                  "be input or ext_oneapi_source",
-                                  PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Internal error. The target state should not be input "
+                        "or ext_oneapi_source");
         break;
       }
     }
diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp
index 6783d2f1ef8e4..d5934472fbefa 100644
--- a/sycl/source/detail/memory_manager.cpp
+++ b/sycl/source/detail/memory_manager.cpp
@@ -459,7 +459,9 @@ void *MemoryManager::allocateMemSubBuffer(ContextImplPtr TargetContext,
         Error);
 
   if (Error != PI_SUCCESS) {
-    Plugin->reportPiError(Error, "allocateMemSubBuffer()");
+    throw set_pi_error(exception(make_error_code(errc::runtime),
+                                 "allocateMemSubBuffer() failed"),
+                       Error);
   }
 
   return NewMem;
@@ -750,8 +752,8 @@ static void copyH2H(SYCLMemObjI *, char *SrcMem, QueueImplPtr,
   if ((DimSrc != 1 || DimDst != 1) &&
       (SrcOffset != id<3>{0, 0, 0} || DstOffset != id<3>{0, 0, 0} ||
        SrcSize != SrcAccessRange || DstSize != DstAccessRange)) {
-    throw runtime_error("Not supported configuration of memcpy requested",
-                        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::feature_not_supported),
+                    "Not supported configuration of memcpy requested");
   }
 
   SrcMem += SrcOffset[0] * SrcElemSize;
@@ -842,8 +844,8 @@ void MemoryManager::fill(SYCLMemObjI *SYCLMemObj, void *Mem, QueueImplPtr Queue,
     }
     // The sycl::handler uses a parallel_for kernel in the case of unusable
     // Range or Offset, not CG:Fill. So we should not be here.
-    throw runtime_error("Not supported configuration of fill requested",
-                        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime),
+                    "Not supported configuration of fill requested");
   } else {
     if (OutEventImpl != nullptr)
       OutEventImpl->setHostEnqueueTime();
@@ -863,8 +865,8 @@ void *MemoryManager::map(SYCLMemObjI *, void *Mem, QueueImplPtr Queue,
                          std::vector<sycl::detail::pi::PiEvent> DepEvents,
                          sycl::detail::pi::PiEvent &OutEvent) {
   if (!Queue) {
-    throw runtime_error("Not supported configuration of map requested",
-                        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime),
+                    "Not supported configuration of map requested");
   }
 
   pi_map_flags Flags = 0;
@@ -909,8 +911,8 @@ void MemoryManager::unmap(SYCLMemObjI *, void *Mem, QueueImplPtr Queue,
 
   // Execution on host is not supported here.
   if (!Queue) {
-    throw runtime_error("Not supported configuration of unmap requested",
-                        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime),
+                    "Not supported configuration of unmap requested");
   }
   // All DepEvents are to the same Context.
   // Using the plugin of the Queue.
@@ -939,8 +941,8 @@ void MemoryManager::copy_usm(const void *SrcMem, QueueImplPtr SrcQueue,
   }
 
   if (!SrcMem || !DstMem)
-    throw runtime_error("NULL pointer argument in memory copy operation.",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "NULL pointer argument in memory copy operation.");
 
   const PluginPtr &Plugin = SrcQueue->getPlugin();
   if (OutEventImpl != nullptr)
@@ -968,8 +970,8 @@ void MemoryManager::fill_usm(void *Mem, QueueImplPtr Queue, size_t Length,
   }
 
   if (!Mem)
-    throw runtime_error("NULL pointer argument in memory fill operation.",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "NULL pointer argument in memory fill operation.");
   if (OutEventImpl != nullptr)
     OutEventImpl->setHostEnqueueTime();
   const PluginPtr &Plugin = Queue->getPlugin();
@@ -1551,8 +1553,8 @@ void MemoryManager::ext_oneapi_copy_usm_cmd_buffer(
     void *DstMem, std::vector<sycl::detail::pi::PiExtSyncPoint> Deps,
     sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
   if (!SrcMem || !DstMem)
-    throw runtime_error("NULL pointer argument in memory copy operation.",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "NULL pointer argument in memory copy operation.");
 
   const PluginPtr &Plugin = Context->getPlugin();
   pi_result Result =
@@ -1576,8 +1578,8 @@ void MemoryManager::ext_oneapi_fill_usm_cmd_buffer(
     sycl::detail::pi::PiExtSyncPoint *OutSyncPoint) {
 
   if (!DstMem)
-    throw runtime_error("NULL pointer argument in memory fill operation.",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "NULL pointer argument in memory fill operation.");
 
   const PluginPtr &Plugin = Context->getPlugin();
 
@@ -1619,8 +1621,8 @@ void MemoryManager::ext_oneapi_fill_cmd_buffer(
   }
   // The sycl::handler uses a parallel_for kernel in the case of unusable
   // Range or Offset, not CG:Fill. So we should not be here.
-  throw runtime_error("Not supported configuration of fill requested",
-                      PI_ERROR_INVALID_OPERATION);
+  throw exception(make_error_code(errc::runtime),
+                  "Not supported configuration of fill requested");
 }
 
 void MemoryManager::ext_oneapi_prefetch_usm_cmd_buffer(
diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp
index ee0ef9e26c35d..ea16c37ce07c0 100644
--- a/sycl/source/detail/pi.cpp
+++ b/sycl/source/detail/pi.cpp
@@ -504,8 +504,8 @@ template <backend BE> const PluginPtr &getPlugin() {
       return *Plugin;
     }
 
-  throw runtime_error("pi::getPlugin couldn't find plugin",
-                      PI_ERROR_INVALID_OPERATION);
+  throw exception(make_error_code(errc::runtime),
+                  "pi::getPlugin couldn't find plugin");
 }
 
 template __SYCL_EXPORT const PluginPtr &getPlugin<backend::opencl>();
diff --git a/sycl/source/detail/platform_util.cpp b/sycl/source/detail/platform_util.cpp
index 09250ca13d7cf..14fa713674e3e 100644
--- a/sycl/source/detail/platform_util.cpp
+++ b/sycl/source/detail/platform_util.cpp
@@ -40,9 +40,8 @@ static void cpuid(uint32_t *CPUInfo, uint32_t Type, uint32_t SubType = 0) {
 #endif
 
 uint32_t PlatformUtil::getMaxClockFrequency() {
-  throw runtime_error(
-      "max_clock_frequency parameter is not supported for host device",
-      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "max_clock_frequency parameter is not supported on host");
   return 0;
 }
 
diff --git a/sycl/source/detail/plugin.hpp b/sycl/source/detail/plugin.hpp
index d36c9769bb8ba..fa01cb35449a6 100644
--- a/sycl/source/detail/plugin.hpp
+++ b/sycl/source/detail/plugin.hpp
@@ -189,16 +189,6 @@ class plugin {
     __SYCL_CHECK_CODE_THROW_VIA_ERRC(pi_result, errc);
   }
 
-  void reportPiError(sycl::detail::pi::PiResult pi_result,
-                     const char *context) const {
-    if (pi_result != PI_SUCCESS) {
-      throw sycl::runtime_error(std::string(context) +
-                                    " API failed with error: " +
-                                    sycl::detail::codeToString(pi_result),
-                                pi_result);
-    }
-  }
-
   /// Calls the PiApi, traces the call, and returns the result.
   ///
   /// Usage:
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index 7aa27f028b384..94e3e062a0d83 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -93,7 +93,10 @@ createBinaryProgram(const ContextImplPtr Context, const device &Device,
       Metadata.size(), Metadata.data(), &BinaryStatus, &Program);
 
   if (BinaryStatus != CL_SUCCESS) {
-    throw runtime_error("Creating program with binary failed.", BinaryStatus);
+    throw detail::set_pi_error(
+        exception(make_error_code(errc::runtime),
+                  "Creating program with binary failed."),
+        BinaryStatus);
   }
 
   return Program;
@@ -180,12 +183,12 @@ ProgramManager::createPIProgram(const RTDeviceBinaryImage &Img,
 
   // perform minimal sanity checks on the device image and the descriptor
   if (RawImg.BinaryEnd < RawImg.BinaryStart) {
-    throw runtime_error("Malformed device program image descriptor",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::runtime),
+                    "Malformed device program image descriptor");
   }
   if (RawImg.BinaryEnd == RawImg.BinaryStart) {
-    throw runtime_error("Invalid device program image: size is zero",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::runtime),
+                    "Invalid device program image: size is zero");
   }
   size_t ImgSize = Img.getSize();
 
@@ -1074,9 +1077,9 @@ ProgramManager::ProgramManager() : m_AsanFoundInImage(false) {
     std::ifstream File(SpvFile, std::ios::binary);
 
     if (!File.is_open())
-      throw runtime_error(std::string("Can't open file specified via ") +
-                              UseSpvEnv + ": " + SpvFile,
-                          PI_ERROR_INVALID_VALUE);
+      throw exception(make_error_code(errc::runtime),
+                      std::string("Can't open file specified via ") +
+                          UseSpvEnv + ": " + SpvFile);
     File.seekg(0, std::ios::end);
     size_t Size = File.tellg();
     std::unique_ptr<char[]> Data(new char[Size]);
@@ -1084,9 +1087,9 @@ ProgramManager::ProgramManager() : m_AsanFoundInImage(false) {
     File.read(Data.get(), Size);
     File.close();
     if (!File.good())
-      throw runtime_error(std::string("read from ") + SpvFile +
-                              std::string(" failed"),
-                          PI_ERROR_INVALID_VALUE);
+      throw exception(make_error_code(errc::runtime),
+                      std::string("read from ") + SpvFile +
+                          std::string(" failed"));
     // No need for a mutex here since all access to these private fields is
     // blocked until the construction of the ProgramManager singleton is
     // finished.
@@ -1187,8 +1190,8 @@ ProgramManager::getDeviceImage(const std::string &KernelName,
     return *Img;
   }
 
-  throw runtime_error("No kernel named " + KernelName + " was found",
-                      PI_ERROR_INVALID_KERNEL_NAME);
+  throw exception(make_error_code(errc::runtime),
+                  "No kernel named " + KernelName + " was found");
 }
 
 RTDeviceBinaryImage &ProgramManager::getDeviceImage(
@@ -1610,7 +1613,7 @@ void ProgramManager::dumpImage(const RTDeviceBinaryImage &Img,
   std::ofstream F(Fname, std::ios::binary);
 
   if (!F.is_open()) {
-    throw runtime_error("Can not write " + Fname, PI_ERROR_UNKNOWN);
+    throw exception(make_error_code(errc::runtime), "Can not write " + Fname);
   }
   Img.dump(F);
   F.close();
@@ -1693,8 +1696,9 @@ static bool compatibleWithDevice(RTDeviceBinaryImage *BinImage,
           PIDeviceHandle, &DevBin,
           /*num bin images = */ (pi_uint32)1, &SuitableImageID);
   if (Error != PI_SUCCESS && Error != PI_ERROR_INVALID_BINARY)
-    throw runtime_error("Invalid binary image or device",
-                        PI_ERROR_INVALID_VALUE);
+    throw detail::set_pi_error(exception(make_error_code(errc::runtime),
+                                         "Invalid binary image or device"),
+                               Error);
 
   return (0 == SuitableImageID);
 }
@@ -1704,8 +1708,8 @@ kernel_id ProgramManager::getSYCLKernelID(const std::string &KernelName) {
 
   auto KernelID = m_KernelName2KernelIDs.find(KernelName);
   if (KernelID == m_KernelName2KernelIDs.end())
-    throw runtime_error("No kernel found with the specified name",
-                        PI_ERROR_INVALID_KERNEL_NAME);
+    throw exception(make_error_code(errc::runtime),
+                    "No kernel found with the specified name");
 
   return KernelID->second;
 }
@@ -2134,10 +2138,12 @@ ProgramManager::compile(const device_image_plain &DeviceImage,
   if (InputImpl->get_bin_image_ref()->getFormat() !=
           PI_DEVICE_BINARY_TYPE_SPIRV &&
       Devs.size() > 1)
-    sycl::runtime_error(
-        "Creating a program from AOT binary for multiple device is not "
-        "supported",
-        PI_ERROR_INVALID_OPERATION);
+    // FIXME: It was probably intended to be thrown, but a unittest starts
+    // failing if we do so, investigate independently of switching to SYCL 2020
+    // `exception`.
+    exception(make_error_code(errc::feature_not_supported),
+              "Creating a program from AOT binary for multiple device is "
+              "not supported");
 
   // Device is not used when creating program from SPIRV, so passing only one
   // device is OK.
@@ -2230,7 +2236,8 @@ ProgramManager::link(const device_image_plain &DeviceImage,
       const std::string ErrorMsg = getProgramBuildLog(LinkedProg, ContextImpl);
       throw sycl::exception(make_error_code(errc::build), ErrorMsg);
     }
-    Plugin->reportPiError(Error, "link()");
+    throw set_pi_error(exception(make_error_code(errc::build), "link() failed"),
+                       Error);
   }
 
   std::shared_ptr<std::vector<kernel_id>> KernelIDs{new std::vector<kernel_id>};
@@ -2329,10 +2336,12 @@ device_image_plain ProgramManager::build(const device_image_plain &DeviceImage,
     if (InputImpl->get_bin_image_ref()->getFormat() !=
             PI_DEVICE_BINARY_TYPE_SPIRV &&
         Devs.size() > 1)
-      sycl::runtime_error(
-          "Creating a program from AOT binary for multiple device is not "
-          "supported",
-          PI_ERROR_INVALID_OPERATION);
+      // FIXME: It was probably intended to be thrown, but a unittest starts
+      // failing if we do so, investigate independently of switching to SYCL
+      // 2020 `exception`.
+      exception(make_error_code(errc::feature_not_supported),
+                "Creating a program from AOT binary for multiple device "
+                "is not supported");
 
     // Device is not used when creating program from SPIRV, so passing only one
     // device is OK.
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
index 4adb0f2dccac1..ebfa91bc8a19c 100644
--- a/sycl/source/detail/queue_impl.cpp
+++ b/sycl/source/detail/queue_impl.cpp
@@ -222,8 +222,8 @@ event queue_impl::memcpy(const std::shared_ptr<detail::queue_impl> &Self,
 
   if ((!Src || !Dest) && Count != 0) {
     report(CodeLoc);
-    throw runtime_error("NULL pointer argument in memory copy operation.",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "NULL pointer argument in memory copy operation.");
   }
   return submitMemOpHelper(
       Self, DepEvents, CallerNeedsEvent,
diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp
index aead8ad77e1a4..e8ec2d9764afe 100644
--- a/sycl/source/detail/scheduler/commands.cpp
+++ b/sycl/source/detail/scheduler/commands.cpp
@@ -2906,8 +2906,8 @@ pi_int32 ExecCGCommand::enqueueImpCommandBuffer() {
   }
 
   default:
-    throw runtime_error("CG type not implemented for command buffers.",
-                        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime),
+                    "CG type not implemented for command buffers.");
   }
 }
 
diff --git a/sycl/source/detail/scheduler/graph_processor.cpp b/sycl/source/detail/scheduler/graph_processor.cpp
index 9853b25d9d310..18ea174ecd230 100644
--- a/sycl/source/detail/scheduler/graph_processor.cpp
+++ b/sycl/source/detail/scheduler/graph_processor.cpp
@@ -36,7 +36,7 @@ void Scheduler::GraphProcessor::waitForEvent(const EventImplPtr &Event,
       enqueueCommand(Cmd, GraphReadLock, Res, ToCleanUp, Cmd, BLOCKING);
   if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
     // TODO: Reschedule commands.
-    throw runtime_error("Enqueue process failed.", PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime), "Enqueue process failed.");
 
   assert(Cmd->getEvent() == Event);
 
diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp
index fbea6f14dea3d..258ab3f6d6a54 100644
--- a/sycl/source/detail/scheduler/scheduler.cpp
+++ b/sycl/source/detail/scheduler/scheduler.cpp
@@ -54,8 +54,8 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record,
     bool Enqueued =
         GraphProcessor::enqueueCommand(Cmd, GraphReadLock, Res, ToCleanUp, Cmd);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.",
-                          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime),
+                      "Enqueue process failed.");
 #ifdef XPTI_ENABLE_INSTRUMENTATION
     // Capture the dependencies
     DepCommands.insert(Cmd);
@@ -67,8 +67,8 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record,
     bool Enqueued =
         GraphProcessor::enqueueCommand(Cmd, GraphReadLock, Res, ToCleanUp, Cmd);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.",
-                          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime),
+                      "Enqueue process failed.");
 #ifdef XPTI_ENABLE_INSTRUMENTATION
     DepCommands.insert(Cmd);
 #endif
@@ -80,8 +80,8 @@ void Scheduler::waitForRecordToFinish(MemObjRecord *Record,
     bool Enqueued = GraphProcessor::enqueueCommand(ReleaseCmd, GraphReadLock,
                                                    Res, ToCleanUp, ReleaseCmd);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.",
-                          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime),
+                      "Enqueue process failed.");
 #ifdef XPTI_ENABLE_INSTRUMENTATION
     // Report these dependencies to the Command so these dependencies can be
     // reported as edges
@@ -171,8 +171,8 @@ void Scheduler::enqueueCommandForCG(EventImplPtr NewEvent,
                                                 Blocking);
       try {
         if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-          throw runtime_error("Auxiliary enqueue process failed.",
-                              PI_ERROR_INVALID_OPERATION);
+          throw exception(make_error_code(errc::runtime),
+                          "Auxiliary enqueue process failed.");
       } catch (...) {
         // enqueueCommand() func and if statement above may throw an exception,
         // so destroy required resources to avoid memory leak
@@ -188,8 +188,8 @@ void Scheduler::enqueueCommandForCG(EventImplPtr NewEvent,
         bool Enqueued = GraphProcessor::enqueueCommand(
             NewCmd, Lock, Res, ToCleanUp, NewCmd, Blocking);
         if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-          throw runtime_error("Enqueue process failed.",
-                              PI_ERROR_INVALID_OPERATION);
+          throw exception(make_error_code(errc::runtime),
+                          "Enqueue process failed.");
       } catch (...) {
         // enqueueCommand() func and if statement above may throw an exception,
         // so destroy required resources to avoid memory leak
@@ -222,15 +222,15 @@ EventImplPtr Scheduler::addCopyBack(Requirement *Req) {
     for (Command *Cmd : AuxiliaryCmds) {
       Enqueued = GraphProcessor::enqueueCommand(Cmd, Lock, Res, ToCleanUp, Cmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
 
     Enqueued =
         GraphProcessor::enqueueCommand(NewCmd, Lock, Res, ToCleanUp, NewCmd);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.",
-                          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime),
+                      "Enqueue process failed.");
   } catch (...) {
     auto WorkerQueue = NewCmd->getEvent()->getWorkerQueue();
     assert(WorkerQueue && "WorkerQueue for CopyBack command must be not null");
@@ -309,16 +309,16 @@ EventImplPtr Scheduler::addHostAccessor(Requirement *Req) {
     for (Command *Cmd : AuxiliaryCmds) {
       Enqueued = GraphProcessor::enqueueCommand(Cmd, Lock, Res, ToCleanUp, Cmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
 
     if (Command *NewCmd = static_cast<Command *>(NewCmdEvent->getCommand())) {
       Enqueued =
           GraphProcessor::enqueueCommand(NewCmd, Lock, Res, ToCleanUp, NewCmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
   }
 
@@ -352,8 +352,8 @@ void Scheduler::enqueueLeavesOfReqUnlocked(const Requirement *const Req,
       bool Enqueued = GraphProcessor::enqueueCommand(Cmd, GraphReadLock, Res,
                                                      ToCleanUp, Cmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
   };
 
@@ -372,8 +372,8 @@ void Scheduler::enqueueUnblockedCommands(
     bool Enqueued =
         GraphProcessor::enqueueCommand(Cmd, GraphReadLock, Res, ToCleanUp, Cmd);
     if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-      throw runtime_error("Enqueue process failed.",
-                          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::runtime),
+                      "Enqueue process failed.");
   }
 }
 
@@ -674,16 +674,16 @@ EventImplPtr Scheduler::addCommandGraphUpdate(
     for (Command *Cmd : AuxiliaryCmds) {
       Enqueued = GraphProcessor::enqueueCommand(Cmd, Lock, Res, ToCleanUp, Cmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
 
     if (Command *NewCmd = static_cast<Command *>(NewCmdEvent->getCommand())) {
       Enqueued =
           GraphProcessor::enqueueCommand(NewCmd, Lock, Res, ToCleanUp, NewCmd);
       if (!Enqueued && EnqueueResultT::SyclEnqueueFailed == Res.MResult)
-        throw runtime_error("Enqueue process failed.",
-                            PI_ERROR_INVALID_OPERATION);
+        throw exception(make_error_code(errc::runtime),
+                        "Enqueue process failed.");
     }
   }
 
diff --git a/sycl/source/detail/spec_constant_impl.cpp b/sycl/source/detail/spec_constant_impl.cpp
index 13fbf9fcc9bdd..43306382a2dea 100644
--- a/sycl/source/detail/spec_constant_impl.cpp
+++ b/sycl/source/detail/spec_constant_impl.cpp
@@ -22,8 +22,8 @@ namespace detail {
 
 void spec_constant_impl::set(size_t Size, const void *Val) {
   if (0 == Size)
-    throw sycl::runtime_error("invalid spec constant size",
-                              PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "invalid spec constant size");
   auto *BytePtr = reinterpret_cast<const char *>(Val);
   this->Bytes.assign(BytePtr, BytePtr + Size);
 }
diff --git a/sycl/source/detail/sycl_mem_obj_t.hpp b/sycl/source/detail/sycl_mem_obj_t.hpp
index 88cda02210dd7..655ab6b24e8b1 100644
--- a/sycl/source/detail/sycl_mem_obj_t.hpp
+++ b/sycl/source/detail/sycl_mem_obj_t.hpp
@@ -249,10 +249,9 @@ class SYCLMemObjT : public SYCLMemObjI {
     MHostPtrReadOnly = IsConstPtr;
     setAlign(RequiredAlign);
     if (useHostPtr())
-      throw runtime_error(
-          "Buffer constructor from a pair of iterator values does not support "
-          "use_host_ptr property.",
-          PI_ERROR_INVALID_OPERATION);
+      throw exception(make_error_code(errc::invalid),
+                      "Buffer constructor from a pair of iterator values does "
+                      "not support use_host_ptr property.");
 
     setAlign(RequiredAlign);
     MShadowCopy = allocateHostMem();
@@ -277,7 +276,7 @@ class SYCLMemObjT : public SYCLMemObjI {
     (void)InitFromUserData;
     (void)HostPtr;
     (void)InteropEvent;
-    throw runtime_error("Not implemented", PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime), "Not implemented");
   }
 
   MemObjType getType() const override { return MemObjType::Undefined; }
diff --git a/sycl/source/detail/usm/usm_impl.cpp b/sycl/source/detail/usm/usm_impl.cpp
index 7237e88be440f..5039ce78f21cd 100644
--- a/sycl/source/detail/usm/usm_impl.cpp
+++ b/sycl/source/detail/usm/usm_impl.cpp
@@ -554,7 +554,9 @@ alloc get_pointer_type(const void *Ptr, const context &Ctxt) {
     return alloc::unknown;
   // otherwise PI_SUCCESS is expected
   if (Err != PI_SUCCESS) {
-    Plugin->reportPiError(Err, "get_pointer_type()");
+    throw detail::set_pi_error(
+        exception(make_error_code(errc::runtime), "get_pointer_type() failed"),
+        Err);
   }
 
   alloc ResultAlloc;
@@ -583,8 +585,8 @@ 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!",
-                        PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "Ptr not a valid USM allocation!");
 
   std::shared_ptr<detail::context_impl> CtxImpl = detail::getSyclObjImpl(Ctxt);
 
@@ -592,8 +594,8 @@ 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!",
-                          PI_ERROR_INVALID_VALUE);
+      throw exception(make_error_code(errc::invalid),
+                      "No devices in passed context!");
 
     // Just return the first device in the context
     return Devs[0];
@@ -614,8 +616,8 @@ device get_pointer_device(const void *Ptr, const context &Ctxt) {
       PltImpl->getDeviceImpl(DeviceId);
   if (DevImpl)
     return detail::createSyclObjFromImpl<device>(DevImpl);
-  throw runtime_error("Cannot find device associated with USM allocation!",
-                      PI_ERROR_INVALID_OPERATION);
+  throw exception(make_error_code(errc::runtime),
+                  "Cannot find device associated with USM allocation!");
 }
 
 // Device copy enhancement APIs, prepare_for and release_from USM.
diff --git a/sycl/source/device_selector.cpp b/sycl/source/device_selector.cpp
index c69910edc2841..e0e4f535232a5 100644
--- a/sycl/source/device_selector.cpp
+++ b/sycl/source/device_selector.cpp
@@ -128,7 +128,7 @@ device select_device(DSelectorInvocableType DeviceSelectorInvocable,
     Message += Acc;
   }
   Message += Suffix;
-  throw sycl::runtime_error(Message, PI_ERROR_DEVICE_NOT_FOUND);
+  throw exception(make_error_code(errc::runtime), Message);
 }
 
 // select_device(selector)
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index d82f24c9ad661..546ef72a4980b 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -496,9 +496,8 @@ event handler::finalize() {
   }
 
   if (!CommandGroup)
-    throw sycl::runtime_error(
-        "Internal Error. Command group cannot be constructed.",
-        PI_ERROR_INVALID_OPERATION);
+    throw exception(make_error_code(errc::runtime),
+                    "Internal Error. Command group cannot be constructed.");
 
   // If there is a graph associated with the handler we are in the explicit
   // graph mode, so we store the CG instead of submitting it to the scheduler,
@@ -788,7 +787,8 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
     break;
   }
   case kernel_param_kind_t::kind_invalid:
-    throw runtime_error("Invalid kernel param kind", PI_ERROR_INVALID_VALUE);
+    throw exception(make_error_code(errc::invalid),
+                    "Invalid kernel param kind");
     break;
   }
 }
diff --git a/sycl/test-e2e/Config/env_vars.cpp b/sycl/test-e2e/Config/env_vars.cpp
index 5a892965064df..3c2c290aff917 100644
--- a/sycl/test-e2e/Config/env_vars.cpp
+++ b/sycl/test-e2e/Config/env_vars.cpp
@@ -23,25 +23,17 @@ int main() {
   int data = 5;
   buffer<int, 1> buf(&data, range<1>(1));
   queue myQueue;
-  if (getenv("SHOULD_CRASH")) {
-    try {
-      myQueue.submit([&](handler &cgh) {
-        auto B = buf.get_access<access::mode::read_write>(cgh);
-        cgh.single_task<class kernel_name1>([=]() { B[0] = 0; });
-      });
-    } catch (sycl::runtime_error &e) {
-      // Exit immediately, otherwise the buffer destructor may actually try to
-      // enqueue the command once again, and throw another exception.
-      exit(0);
-    } catch (sycl::exception &e) {
-      if (e.code() == errc::build)
-        exit(0);
-    }
-    assert(0 && "Expected exception was *not* thrown");
-  } else {
+  bool shouldCrash = getenv("SHOULD_CRASH");
+  try {
     myQueue.submit([&](handler &cgh) {
       auto B = buf.get_access<access::mode::read_write>(cgh);
-      cgh.single_task<class kernel_name2>([=]() { B[0] = 0; });
+      cgh.single_task<class kernel_name1>([=]() { B[0] = 0; });
     });
+    assert(!shouldCrash);
+  } catch (sycl::exception &e) {
+    assert(shouldCrash);
+    assert(e.code() == errc::build);
   }
+
+  return 0;
 }
diff --git a/sycl/test-e2e/FilterSelector/select.cpp b/sycl/test-e2e/FilterSelector/select.cpp
index 106b587afc000..bccd564e89a04 100644
--- a/sycl/test-e2e/FilterSelector/select.cpp
+++ b/sycl/test-e2e/FilterSelector/select.cpp
@@ -116,7 +116,7 @@ int main() {
     // pick something crazy
     device d10(filter_selector("bob:gpu"));
   } catch (const sycl::exception &e) {
-    assert(e.code() == sycl::errc::runtime);
+    assert(e.code() == sycl::errc::invalid);
     const char *ErrorMesg = "Invalid filter string!";
     assert(std::string{e.what()}.find(ErrorMesg) == 0 &&
            "filter_selector(\"bob:gpu\") unexpectedly selected a device");
@@ -126,7 +126,7 @@ int main() {
     // pick something crazy
     device d11(filter_selector("opencl:bob"));
   } catch (const sycl::exception &e) {
-    assert(e.code() == sycl::errc::runtime);
+    assert(e.code() == sycl::errc::invalid);
     const char *ErrorMesg = "Invalid filter string!";
     assert(std::string{e.what()}.find(ErrorMesg) == 0 &&
            "filter_selector(\"opencl:bob\") unexpectedly selected a device");
diff --git a/sycl/test-e2e/Graph/RecordReplay/host_task_last.cpp b/sycl/test-e2e/Graph/RecordReplay/host_task_last.cpp
index 921acca503bf7..0f0ba85639cbf 100644
--- a/sycl/test-e2e/Graph/RecordReplay/host_task_last.cpp
+++ b/sycl/test-e2e/Graph/RecordReplay/host_task_last.cpp
@@ -2,8 +2,10 @@
 // RUN: %{run} %t.out
 // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+
+// Disabled due to https://github.com/intel/llvm/issues/14473
 // Extra run to check for immediate-command-list in Level Zero
-// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// xRUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
 // REQUIRES: aspect-usm_shared_allocations
 
diff --git a/sycl/test-e2e/Tracing/code_location_queue_submit.cpp b/sycl/test-e2e/Tracing/code_location_queue_submit.cpp
index ce780f5e81725..3a018604df1f9 100644
--- a/sycl/test-e2e/Tracing/code_location_queue_submit.cpp
+++ b/sycl/test-e2e/Tracing/code_location_queue_submit.cpp
@@ -24,7 +24,7 @@ int main() {
     // CHECK-DAG:        sycl_device_name : {{.*}}
     // CHECK-DAG:        sycl_context : {{.*}}
     // CHECK: [SYCL] Runtime reports:
-    // CHECK-NEXT: what:  NULL pointer argument in memory copy operation. -30 (PI_ERROR_INVALID_VALUE)
+    // CHECK-NEXT: what:  NULL pointer argument in memory copy operation.
     // CHECK-NEXT: where:{{.*}}code_location_queue_submit.cpp:[[# @LINE + 2 ]] main
     try {
       Q.submit(
diff --git a/sycl/unittests/SYCL2020/KernelID.cpp b/sycl/unittests/SYCL2020/KernelID.cpp
index 1791205234c28..1f493404f367f 100644
--- a/sycl/unittests/SYCL2020/KernelID.cpp
+++ b/sycl/unittests/SYCL2020/KernelID.cpp
@@ -261,9 +261,7 @@ TEST(KernelID, GetKernelIDInvalidKernelName) {
     FAIL() << "Expected an exception";
   } catch (sycl::exception const &e) {
     EXPECT_TRUE(e.code() == sycl::errc::runtime);
-    EXPECT_EQ(std::string("No kernel found with the specified name -46 "
-                          "(PI_ERROR_INVALID_KERNEL_NAME)"),
-              e.what());
+    EXPECT_EQ(std::string("No kernel found with the specified name"), e.what());
   } catch (...) {
     FAIL() << "Expected sycl::exception";
   }