@@ -644,125 +644,6 @@ struct sub_group {
644
644
#endif
645
645
}
646
646
647
- #ifndef __INTEL_PREVIEW_BREAKING_CHANGES__
648
- /* --- deprecated collective functions --- */
649
- template <typename T>
650
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
651
- " sycl::ext::oneapi::broadcast instead." )
652
- EnableIfIsScalarArithmetic<T> broadcast(T x, id<1 > local_id) const {
653
- #ifdef __SYCL_DEVICE_ONLY__
654
- return sycl::detail::spirv::GroupBroadcast<sub_group>(x, local_id);
655
- #else
656
- (void )x;
657
- (void )local_id;
658
- throw sycl::exception (make_error_code (errc::feature_not_supported),
659
- " Sub-groups are not supported on host." );
660
- #endif
661
- }
662
-
663
- template <typename T, class BinaryOperation >
664
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
665
- " sycl::ext::oneapi::reduce instead." )
666
- EnableIfIsScalarArithmetic<T> reduce(T x, BinaryOperation op) const {
667
- #ifdef __SYCL_DEVICE_ONLY__
668
- return sycl::detail::calc<__spv::GroupOperation::Reduce>(
669
- typename sycl::detail::GroupOpTag<T>::type (), *this , x, op);
670
- #else
671
- (void )x;
672
- (void )op;
673
- throw sycl::exception (make_error_code (errc::feature_not_supported),
674
- " Sub-groups are not supported on host." );
675
- #endif
676
- }
677
-
678
- template <typename T, class BinaryOperation >
679
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
680
- " sycl::ext::oneapi::reduce instead." )
681
- EnableIfIsScalarArithmetic<T> reduce(T x, T init, BinaryOperation op) const {
682
- #ifdef __SYCL_DEVICE_ONLY__
683
- return op (init, reduce (x, op));
684
- #else
685
- (void )x;
686
- (void )init;
687
- (void )op;
688
- throw sycl::exception (make_error_code (errc::feature_not_supported),
689
- " Sub-groups are not supported on host." );
690
- #endif
691
- }
692
-
693
- template <typename T, class BinaryOperation >
694
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
695
- " sycl::ext::oneapi::exclusive_scan instead." )
696
- EnableIfIsScalarArithmetic<T> exclusive_scan(T x, BinaryOperation op) const {
697
- #ifdef __SYCL_DEVICE_ONLY__
698
- return sycl::detail::calc<__spv::GroupOperation::ExclusiveScan>(
699
- typename sycl::detail::GroupOpTag<T>::type (), *this , x, op);
700
- #else
701
- (void )x;
702
- (void )op;
703
- throw sycl::exception (make_error_code (errc::feature_not_supported),
704
- " Sub-groups are not supported on host." );
705
- #endif
706
- }
707
-
708
- template <typename T, class BinaryOperation >
709
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
710
- " sycl::ext::oneapi::exclusive_scan instead." )
711
- EnableIfIsScalarArithmetic<T> exclusive_scan(T x, T init,
712
- BinaryOperation op) const {
713
- #ifdef __SYCL_DEVICE_ONLY__
714
- if (get_local_id ().get (0 ) == 0 ) {
715
- x = op (init, x);
716
- }
717
- T scan = exclusive_scan (x, op);
718
- if (get_local_id ().get (0 ) == 0 ) {
719
- scan = init;
720
- }
721
- return scan;
722
- #else
723
- (void )x;
724
- (void )init;
725
- (void )op;
726
- throw sycl::exception (make_error_code (errc::feature_not_supported),
727
- " Sub-groups are not supported on host." );
728
- #endif
729
- }
730
-
731
- template <typename T, class BinaryOperation >
732
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
733
- " sycl::ext::oneapi::inclusive_scan instead." )
734
- EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op) const {
735
- #ifdef __SYCL_DEVICE_ONLY__
736
- return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
737
- typename sycl::detail::GroupOpTag<T>::type (), *this , x, op);
738
- #else
739
- (void )x;
740
- (void )op;
741
- throw sycl::exception (make_error_code (errc::feature_not_supported),
742
- " Sub-groups are not supported on host." );
743
- #endif
744
- }
745
-
746
- template <typename T, class BinaryOperation >
747
- __SYCL_DEPRECATED (" Collectives in the sub-group class are deprecated. Use "
748
- " sycl::ext::oneapi::inclusive_scan instead." )
749
- EnableIfIsScalarArithmetic<T> inclusive_scan(T x, BinaryOperation op,
750
- T init) const {
751
- #ifdef __SYCL_DEVICE_ONLY__
752
- if (get_local_id ().get (0 ) == 0 ) {
753
- x = op (init, x);
754
- }
755
- return inclusive_scan (x, op);
756
- #else
757
- (void )x;
758
- (void )op;
759
- (void )init;
760
- throw sycl::exception (make_error_code (errc::feature_not_supported),
761
- " Sub-groups are not supported on host." );
762
- #endif
763
- }
764
- #endif // __INTEL_PREVIEW_BREAKING_CHANGES__
765
-
766
647
linear_id_type get_group_linear_range () const {
767
648
#ifdef __SYCL_DEVICE_ONLY__
768
649
return static_cast <linear_id_type>(get_group_range ()[0 ]);
0 commit comments