diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp index 908b5d4aebff0..9781195d10c4c 100644 --- a/sycl/include/sycl/sub_group.hpp +++ b/sycl/include/sycl/sub_group.hpp @@ -219,8 +219,9 @@ struct sub_group { /* --- one-input shuffles --- */ /* indices in [0 , sub_group size) */ - - template T shuffle(T x, id_type local_id) const { + template + __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.") + T shuffle(T x, id_type local_id) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::SubgroupShuffle(x, local_id); #else @@ -231,7 +232,9 @@ struct sub_group { #endif } - template T shuffle_down(T x, uint32_t delta) const { + template + __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.") + T shuffle_down(T x, uint32_t delta) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::SubgroupShuffleDown(x, delta); #else @@ -242,7 +245,9 @@ struct sub_group { #endif } - template T shuffle_up(T x, uint32_t delta) const { + template + __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.") + T shuffle_up(T x, uint32_t delta) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::SubgroupShuffleUp(x, delta); #else @@ -253,7 +258,9 @@ struct sub_group { #endif } - template T shuffle_xor(T x, id_type value) const { + template + __SYCL_DEPRECATED("Shuffles in the sub-group class are deprecated.") + T shuffle_xor(T x, id_type value) const { #ifdef __SYCL_DEVICE_ONLY__ return sycl::detail::spirv::SubgroupShuffleXor(x, value); #else @@ -358,7 +365,7 @@ struct sub_group { } return res; } -#else // __NVPTX__ || __AMDGCN__ +#else // __NVPTX__ || __AMDGCN__ template > std::enable_if_t< @@ -608,6 +615,7 @@ struct sub_group { } /* --- synchronization functions --- */ + __SYCL_DEPRECATED("Sub-group barrier with no arguments is deprecated.") void barrier() const { #ifdef __SYCL_DEVICE_ONLY__ __spirv_ControlBarrier( @@ -636,6 +644,7 @@ struct sub_group { #endif } +#ifndef __INTEL_PREVIEW_BREAKING_CHANGES__ /* --- deprecated collective functions --- */ template __SYCL_DEPRECATED("Collectives in the sub-group class are deprecated. Use " @@ -752,6 +761,7 @@ struct sub_group { "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__