Skip to content

Commit

Permalink
[ABI-Break][SYCL] Remove collectives in the sub-group class (#13665)
Browse files Browse the repository at this point in the history
Re-lands #13199 and reverts
#13464.
  • Loading branch information
aelovikov-intel authored Jun 27, 2024
1 parent b429484 commit 7193c26
Showing 1 changed file with 0 additions and 119 deletions.
119 changes: 0 additions & 119 deletions sycl/include/sycl/sub_group.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -643,125 +643,6 @@ 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]);
Expand Down

0 comments on commit 7193c26

Please sign in to comment.