Revert "[ABI-Break][SYCL] Remove collectives in the sub-group class" (#…

…13464) Reverts #13199
intel · Apr 18, 2024 · cb2eede · cb2eede
1 parent a98c87d
commit cb2eede
Showing 1 changed file with 119 additions and 0 deletions.
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
@@ -644,6 +644,125 @@ struct sub_group {
 #endif
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES__
+  /* --- deprecated collective functions --- */
+  template <typename T>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::broadcast instead.")
+  EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return sycl::detail::spirv::GroupBroadcast<sub_group>(x, local_id);
+#else
+    (void)x;
+    (void)local_id;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::reduce instead.")
+  EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return sycl::detail::calc<__spv::GroupOperation::Reduce>(
+        typename sycl::detail::GroupOpTag<T>::type(), *this, x, op);
+#else
+    (void)x;
+    (void)op;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::reduce instead.")
+  EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return op(init, reduce(x, op));
+#else
+    (void)x;
+    (void)init;
+    (void)op;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::exclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return sycl::detail::calc<__spv::GroupOperation::ExclusiveScan>(
+        typename sycl::detail::GroupOpTag<T>::type(), *this, x, op);
+#else
+    (void)x;
+    (void)op;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::exclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
+                                               BinaryOperation op) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    if (get_local_id().get(0) == 0) {
+      x = op(init, x);
+    }
+    T scan = exclusive_scan(x, op);
+    if (get_local_id().get(0) == 0) {
+      scan = init;
+    }
+    return scan;
+#else
+    (void)x;
+    (void)init;
+    (void)op;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::inclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
+        typename sycl::detail::GroupOpTag<T>::type(), *this, x, op);
+#else
+    (void)x;
+    (void)op;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::ext::oneapi::inclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
+                                               T init) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    if (get_local_id().get(0) == 0) {
+      x = op(init, x);
+    }
+    return inclusive_scan(x, op);
+#else
+    (void)x;
+    (void)op;
+    (void)init;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES__
+
   linear_id_type get_group_linear_range() const {
 #ifdef __SYCL_DEVICE_ONLY__
     return static_cast<linear_id_type>(get_group_range()[0]);