[SYCL] Restore sub-group collectives

Signed-off-by: John Pennycook <[email protected]>

Author: Pennycook

sycl/include/CL/sycl/intel/sub_group.hpp

@@ -470,6 +470,74 @@ struct sub_group {
 #endif
   }
 
+  /* --- deprecated collective functions --- */
+  template <typename T>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::broadcast instead.")
+  EnableIfIsScalarArithmetic<T> broadcast(T x, id<1> local_id) const {
+    return sycl::detail::spirv::GroupBroadcast<sub_group>(x, local_id);
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::reduce instead.")
+  EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
+    return sycl::detail::calc<T, __spv::GroupOperation::Reduce,
+                              __spv::Scope::Subgroup>(
+        typename sycl::detail::GroupOpTag<T>::type(), x, op);
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::reduce instead.")
+  EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
+    return op(init, reduce(x, op));
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::exclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
+    return sycl::detail::calc<T, __spv::GroupOperation::ExclusiveScan,
+                              __spv::Scope::Subgroup>(
+        typename sycl::detail::GroupOpTag<T>::type(), x, op);
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::exclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
+                                               BinaryOperation op) const {
+    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;
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::inclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
+    return sycl::detail::calc<T, __spv::GroupOperation::InclusiveScan,
+                              __spv::Scope::Subgroup>(
+        typename sycl::detail::GroupOpTag<T>::type(), x, op);
+  }
+
+  template <typename T, class BinaryOperation>
+  __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use "
+                    "sycl::intel::inclusive_scan instead.")
+  EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
+                                               T init) const {
+    if (get_local_id().get(0) == 0) {
+      x = op(init, x);
+    }
+    return inclusive_scan(x, op);
+  }
+
 protected:
   template <int dimensions> friend class cl::sycl::nd_item;
   sub_group() = default;