diff --git a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
index 1a24e11d2da26..7a6f51cacc39e 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/ballot_group.hpp
@@ -9,10 +9,9 @@
 #pragma once
 
 #include <sycl/aspects.hpp>
-#include <sycl/detail/pi.h> // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
-#include <sycl/exception.hpp>          // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for GetMask
 #include <sycl/ext/oneapi/sub_group_mask.hpp> // for sub_group_mask
 #include <sycl/id.hpp>                        // for id
@@ -49,8 +48,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return (Predicate) ? 1 : 0;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -58,8 +57,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::CallerPositionInMask(Mask);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -67,8 +66,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return 2;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -76,8 +75,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return Mask.count();
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -85,8 +84,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -94,8 +93,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -103,8 +102,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -112,8 +111,8 @@ template <typename ParentGroup> class ballot_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -122,8 +121,8 @@ template <typename ParentGroup> class ballot_group {
     uint32_t Lowest = static_cast<uint32_t>(Mask.find_low()[0]);
     return __spirv_SubgroupLocalInvocationId() == Lowest;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -165,8 +164,8 @@ get_ballot_group(Group group, bool predicate) {
 #endif
 #else
   (void)predicate;
-  throw runtime_error("Non-uniform groups are not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Non-uniform groups are not supported on host.");
 #endif
 }
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
index 5878a4d63c94a..6ba75098bd534 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp
@@ -71,9 +71,9 @@ namespace ext::oneapi::experimental {
 //
 // - OpenCL spec defines several additional features, like, for example, 'v'
 // modifier which allows to print OpenCL vectors: note that these features are
-// not available on host device and therefore their usage should be either
-// guarded using __SYCL_DEVICE_ONLY__ preprocessor macro or avoided in favor
-// of more portable solutions if needed
+// not available on host and therefore their usage should be either guarded
+// using __SYCL_DEVICE_ONLY__ preprocessor macro or avoided in favor of more
+// portable solutions if needed
 //
 template <typename FormatT, typename... Args>
 int printf(const FormatT *__format, Args... args) {
diff --git a/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
index 2c78e19be78ed..56f30f2091a6b 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/fixed_size_group.hpp
@@ -9,10 +9,9 @@
 #pragma once
 
 #include <sycl/aspects.hpp>
-#include <sycl/detail/pi.h>            // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp> // for is_fixed_size_group, is_group
-#include <sycl/exception.hpp>          // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp> // for sub_group_mask
 #include <sycl/id.hpp>                        // for id
@@ -50,8 +49,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv_SubgroupLocalInvocationId() / PartitionSize;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -59,8 +58,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv_SubgroupLocalInvocationId() % PartitionSize;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -68,8 +67,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return __spirv_SubgroupSize() / PartitionSize;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -77,8 +76,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return PartitionSize;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -86,8 +85,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -95,8 +94,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -104,8 +103,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -113,8 +112,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -122,8 +121,8 @@ template <size_t PartitionSize, typename ParentGroup> class fixed_size_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return get_local_linear_id() == 0;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -168,8 +167,8 @@ get_fixed_size_group(Group group) {
   return fixed_size_group<PartitionSize, sycl::sub_group>();
 #endif
 #else
-  throw runtime_error("Non-uniform groups are not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Non-uniform groups are not supported on host.");
 #endif
 }
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
index d85b9513f070a..54b08a4652d09 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp
@@ -9,10 +9,9 @@
 #pragma once
 
 #include <sycl/aspects.hpp>
-#include <sycl/detail/pi.h>            // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
-#include <sycl/exception.hpp>          // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp> // for this_sub_group
 #include <sycl/ext/oneapi/sub_group_mask.hpp> // for sub_group_mask
@@ -55,8 +54,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<id_type>(0);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -64,8 +63,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::CallerPositionInMask(Mask);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -73,8 +72,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return 1;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -82,8 +81,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return Mask.count();
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -91,8 +90,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -100,8 +99,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -109,8 +108,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -118,8 +117,8 @@ class opportunistic_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -128,8 +127,8 @@ class opportunistic_group {
     uint32_t Lowest = static_cast<uint32_t>(Mask.find_low()[0]);
     return __spirv_SubgroupLocalInvocationId() == Lowest;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -162,8 +161,8 @@ inline opportunistic_group get_opportunistic_group() {
   return opportunistic_group(mask);
 #endif
 #else
-  throw runtime_error("Non-uniform groups are not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Non-uniform groups are not supported on host.");
 #endif
 }
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
index 0e369e643cdcb..558396bb6f9c8 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/root_group.hpp
@@ -115,8 +115,8 @@ void group_barrier(ext::oneapi::experimental::root_group<dimensions> G,
 #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::runtime),
+                        "Barriers are not supported on host");
 #endif
 }
 
diff --git a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
index f0553616fc2e1..59af6bdfc753b 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/tangle_group.hpp
@@ -9,10 +9,9 @@
 #pragma once
 
 #include <sycl/aspects.hpp>
-#include <sycl/detail/pi.h>            // for PI_ERROR_INVALID_DEVICE
 #include <sycl/detail/spirv.hpp>
 #include <sycl/detail/type_traits.hpp> // for is_group, is_user_cons...
-#include <sycl/exception.hpp>          // for runtime_error
+#include <sycl/exception.hpp>
 #include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp> // for sub_group_mask
 #include <sycl/id.hpp>                        // for id
@@ -48,8 +47,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<id_type>(0);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -57,8 +56,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return sycl::detail::CallerPositionInMask(Mask);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -66,8 +65,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return 1;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -75,8 +74,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return Mask.count();
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -84,8 +83,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -93,8 +92,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_id()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -102,8 +101,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -111,8 +110,8 @@ template <typename ParentGroup> class tangle_group {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_local_range()[0]);
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -121,8 +120,8 @@ template <typename ParentGroup> class tangle_group {
     uint32_t Lowest = static_cast<uint32_t>(Mask.find_low()[0]);
     return __spirv_SubgroupLocalInvocationId() == Lowest;
 #else
-    throw runtime_error("Non-uniform groups are not supported on host device.",
-                        PI_ERROR_INVALID_DEVICE);
+    throw exception(make_error_code(errc::runtime),
+                    "Non-uniform groups are not supported on host.");
 #endif
   }
 
@@ -158,8 +157,8 @@ get_tangle_group(Group group) {
   return tangle_group<sycl::sub_group>(0);
 #endif
 #else
-  throw runtime_error("Non-uniform groups are not supported on host device.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Non-uniform groups are not supported on host.");
 #endif
 
 } // namespace this_kernel
diff --git a/sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp b/sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp
index 9d8a9f870fcbd..ce4a8e7d99ace 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp
@@ -37,8 +37,8 @@ T reduce_over_group_impl(GroupHelper group_helper, T x, size_t num_elements,
   std::ignore = x;
   std::ignore = num_elements;
   std::ignore = binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Group algorithms are not supported on host.");
 #endif
 }
 } // namespace detail
@@ -55,8 +55,8 @@ reduce_over_group(GroupHelper group_helper, T x, BinaryOperation binary_op) {
       group_helper, x, group_helper.get_group().get_local_linear_range(),
       binary_op);
 #else
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -74,8 +74,8 @@ reduce_over_group(GroupHelper group_helper, V x, T init,
   return binary_op(init, reduce_over_group(group_helper, x, binary_op));
 #else
   std::ignore = group_helper;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -111,8 +111,8 @@ joint_reduce(GroupHelper group_helper, Ptr first, Ptr last,
   std::ignore = first;
   std::ignore = last;
   std::ignore = binary_op;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Group algorithms are not supported on host.");
 #endif
 }
 
@@ -131,8 +131,8 @@ joint_reduce(GroupHelper group_helper, Ptr first, Ptr last, T init,
 #else
   std::ignore = group_helper;
   std::ignore = last;
-  throw runtime_error("Group algorithms are not supported on host.",
-                      PI_ERROR_INVALID_DEVICE);
+  throw exception(make_error_code(errc::runtime),
+                  "Group algorithms are not supported on host.");
 #endif
 }
 } // namespace ext::oneapi::experimental