diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
index b11908fce90f0..a441e5429d457 100644
--- a/sycl/include/sycl/detail/spirv.hpp
+++ b/sycl/include/sycl/detail/spirv.hpp
@@ -9,7 +9,6 @@
 #pragma once
 
 #ifdef __SYCL_DEVICE_ONLY__
-
 // Some __spirv_* inrinsics are automatically forward-declared by the compiler,
 // but not all of them. For example:
 //   __spirv_AtomicStore(unsigned long long*, ...)
@@ -18,7 +17,10 @@
 #include <sycl/__spirv/spirv_ops.hpp>
 #include <sycl/__spirv/spirv_types.hpp>
 
-#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for IdToMaskPosition
+#include <sycl/access/access.hpp>
+#include <sycl/detail/generic_type_traits.hpp>
+#include <sycl/id.hpp>
+#include <sycl/multi_ptr.hpp>
 
 #if defined(__NVPTX__)
 #include <sycl/ext/oneapi/experimental/cuda/masked_shuffles.hpp>
@@ -33,6 +35,7 @@ struct sub_group;
 namespace ext {
 namespace oneapi {
 struct sub_group;
+struct sub_group_mask;
 namespace experimental {
 template <typename ParentGroup> class fragment;
 
@@ -61,6 +64,9 @@ GetMultiPtrDecoratedAs(multi_ptr<FromT, Space, IsDecorated> MPtr) {
 
 template <typename NonUniformGroup>
 inline uint32_t IdToMaskPosition(NonUniformGroup Group, uint32_t Id);
+template <typename NonUniformGroup>
+inline ext::oneapi::sub_group_mask GetMask(NonUniformGroup Group);
+inline sycl::vec<unsigned, 4> ExtractMask(ext::oneapi::sub_group_mask Mask);
 
 namespace spirv {
 
diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
index 48694b0051295..544e7f7a152ac 100644
--- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
+++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
@@ -9,10 +9,10 @@
 
 #include <sycl/detail/helpers.hpp> // for Builder
 #include <sycl/detail/memcpy.hpp>  // detail::memcpy
-#include <sycl/exception.hpp>      // for errc, exception
-#include <sycl/feature_test.hpp>   // for SYCL_EXT_ONEAPI_SUB_GROUP_MASK
-#include <sycl/id.hpp>             // for id
-#include <sycl/marray.hpp>         // for marray
+#include <sycl/detail/spirv.hpp>
+#include <sycl/feature_test.hpp> // for SYCL_EXT_ONEAPI_SUB_GROUP_MASK
+#include <sycl/id.hpp>           // for id
+#include <sycl/marray.hpp>       // for marray
 #include <sycl/sub_group.hpp>
 #include <sycl/vector.hpp> // for vec
 
@@ -378,19 +378,11 @@ group_ballot([[maybe_unused]] Group g, [[maybe_unused]] bool predicate) {
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::commonGroupBallotImpl(g, predicate);
 #else
-  throw exception{errc::feature_not_supported,
-                  "Sub-group mask is not supported on host device"};
+  // Groups are not user-constructible, this call should not be reachable from
+  // host and therefore we do nothing here.
 #endif
 }
 
 } // namespace ext::oneapi
 } // namespace _V1
 } // namespace sycl
-
-// We have a cyclic dependency with
-//   sub_group_mask.hpp
-//   detail/spirv.hpp
-//   non_uniform_groups.hpp
-// "Break" it by including this at the end (instead of beginning). Ideally, we
-// should refactor this somehow...
-#include <sycl/detail/spirv.hpp>