From eb3962aa65c61c761091e9f44c7bf12c264f6082 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov@intel.com>
Date: Fri, 13 Jun 2025 15:10:38 +0200
Subject: [PATCH 1/2] [SYCL][NFC] Optimize dependencies of atomic_ref

---
 sycl/include/sycl/detail/spirv.hpp            |  6 ++++-
 .../sycl/ext/oneapi/sub_group_mask.hpp        | 26 ++++++-------------
 2 files changed, 13 insertions(+), 19 deletions(-)

diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
index 541668ed374d9..535bd3bd09f07 100644
--- a/sycl/include/sycl/detail/spirv.hpp
+++ b/sycl/include/sycl/detail/spirv.hpp
@@ -10,7 +10,11 @@
 
 #ifdef __SYCL_DEVICE_ONLY__
 
-#include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp> // for IdToMaskPosition
+#include <sycl/__spirv/spirv_types.hpp>
+#include <sycl/access/access.hpp>
+#include <sycl/id.hpp>
+#include <sycl/multi_ptr.hpp>
+#include <sycl/detail/generic_type_traits.hpp>
 
 #if defined(__NVPTX__)
 #include <sycl/ext/oneapi/experimental/cuda/masked_shuffles.hpp>
diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
index 63e56327168e8..e0915e2979cd8 100644
--- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
+++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
@@ -9,11 +9,11 @@
 
 #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/vector.hpp>         // for vec
+#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/vector.hpp>       // for vec
 
 #include <assert.h>     // for assert
 #include <climits>      // for CHAR_BIT
@@ -342,8 +342,7 @@ template <typename Group>
 std::enable_if_t<std::is_same_v<std::decay_t<Group>, sub_group> ||
                      std::is_same_v<std::decay_t<Group>, sycl::sub_group>,
                  sub_group_mask>
-group_ballot(Group g, bool predicate) {
-  (void)g;
+group_ballot([[maybe_unused]] Group g, [[maybe_unused]] bool predicate) {
 #ifdef __SYCL_DEVICE_ONLY__
   auto res = __spirv_GroupNonUniformBallot(
       sycl::detail::spirv::group_scope<Group>::value, predicate);
@@ -353,20 +352,11 @@ group_ballot(Group g, bool predicate) {
   return sycl::detail::Builder::createSubGroupMask<sub_group_mask>(
       val, g.get_max_local_range()[0]);
 #else
-  (void)predicate;
-  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>

From 3b2ec597612e63ff9257339d43290fe9add3bc59 Mon Sep 17 00:00:00 2001
From: Alexey Sachkov <alexey.sachkov@intel.com>
Date: Mon, 16 Jun 2025 14:54:54 +0200
Subject: [PATCH 2/2] Clang-format and missing forward declarations

---
 sycl/include/sycl/detail/spirv.hpp | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp
index 535bd3bd09f07..6294b1eaf8fdb 100644
--- a/sycl/include/sycl/detail/spirv.hpp
+++ b/sycl/include/sycl/detail/spirv.hpp
@@ -12,9 +12,9 @@
 
 #include <sycl/__spirv/spirv_types.hpp>
 #include <sycl/access/access.hpp>
+#include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/id.hpp>
 #include <sycl/multi_ptr.hpp>
-#include <sycl/detail/generic_type_traits.hpp>
 
 #if defined(__NVPTX__)
 #include <sycl/ext/oneapi/experimental/cuda/masked_shuffles.hpp>
@@ -28,6 +28,7 @@ struct sub_group;
 namespace ext {
 namespace oneapi {
 struct sub_group;
+struct sub_group_mask;
 namespace experimental {
 template <typename ParentGroup> class ballot_group;
 template <size_t PartitionSize, typename ParentGroup> class fixed_size_group;
@@ -58,6 +59,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 {