Skip to content

Commit 66ecb82

Browse files
dklochkov-embdklochkov-emb
committed
Revert "[SYCL] fix asserts after logical operation changes"
This reverts commit 118830c.
1 parent e0d9acd commit 66ecb82

File tree

1 file changed

+0
-140
lines changed

1 file changed

+0
-140
lines changed

‎sycl/include/sycl/group_algorithm.hpp‎

Lines changed: 0 additions & 140 deletions
Original file line numberDiff line numberDiff line change
@@ -215,19 +215,9 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
215215
detail::is_native_op<T, BinaryOperation>::value),
216216
T>
217217
reduce_over_group(Group g, T x, BinaryOperation binary_op) {
218-
219-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
220-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
221-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
222-
? std::is_same_v<decltype(binary_op(x, x)), bool>
223-
: std::is_same_v<decltype(binary_op(x, x)), T>,
224-
"Result type of binary_op must match scan accumulation type.");
225-
#else
226218
static_assert(
227219
std::is_same_v<decltype(binary_op(x, x)), T>,
228220
"Result type of binary_op must match reduction accumulation type.");
229-
#endif
230-
231221
#ifdef __SYCL_DEVICE_ONLY__
232222
#if defined(__NVPTX__)
233223
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
@@ -301,18 +291,9 @@ std::enable_if_t<
301291
std::is_convertible_v<V, T>),
302292
T>
303293
reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
304-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
305-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
306-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
307-
? std::is_same_v<decltype(binary_op(init, x)), bool>
308-
: std::is_same_v<decltype(binary_op(init, x)), T>,
309-
"Result type of binary_op must match scan accumulation type.");
310-
#else
311294
static_assert(
312295
std::is_same_v<decltype(binary_op(init, x)), T>,
313296
"Result type of binary_op must match reduction accumulation type.");
314-
#endif
315-
316297
#ifdef __SYCL_DEVICE_ONLY__
317298
return binary_op(init, reduce_over_group(g, T(x), binary_op));
318299
#else
@@ -360,18 +341,9 @@ std::enable_if_t<
360341
detail::is_native_op<T, BinaryOperation>::value),
361342
T>
362343
joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
363-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
364-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
365-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
366-
? std::is_same_v<decltype(binary_op(init, *first)), bool>
367-
: std::is_same_v<decltype(binary_op(init, *first)), T>,
368-
"Result type of binary_op must match scan accumulation type.");
369-
#else
370344
static_assert(
371345
std::is_same_v<decltype(binary_op(init, *first)), T>,
372346
"Result type of binary_op must match reduction accumulation type.");
373-
#endif
374-
375347
#ifdef __SYCL_DEVICE_ONLY__
376348
T partial = detail::identity_for_ga_op<T, BinaryOperation>();
377349
sycl::detail::for_each(
@@ -707,16 +679,8 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
707679
detail::is_native_op<T, BinaryOperation>::value),
708680
T>
709681
exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
710-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
711-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
712-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
713-
? std::is_same_v<decltype(binary_op(x, x)), bool>
714-
: std::is_same_v<decltype(binary_op(x, x)), T>,
715-
"Result type of binary_op must match scan accumulation type.");
716-
#else
717682
static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
718683
"Result type of binary_op must match scan accumulation type.");
719-
#endif
720684
#ifdef __SYCL_DEVICE_ONLY__
721685
#if defined(__NVPTX__)
722686
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
@@ -788,16 +752,8 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
788752
detail::is_native_op<T, BinaryOperation>::value),
789753
T>
790754
exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
791-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
792-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
793-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
794-
? std::is_same_v<decltype(binary_op(x, x)), bool>
795-
: std::is_same_v<decltype(binary_op(x, x)), T>,
796-
"Result type of binary_op must match scan accumulation type.");
797-
#else
798755
static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
799756
"Result type of binary_op must match scan accumulation type.");
800-
#endif
801757
T result;
802758
typename detail::get_scalar_binary_op<BinaryOperation>::type
803759
scalar_binary_op{};
@@ -819,17 +775,8 @@ std::enable_if_t<
819775
std::is_convertible_v<V, T>),
820776
T>
821777
exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
822-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
823-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
824-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
825-
? std::is_same_v<decltype(binary_op(init, x)), bool>
826-
: std::is_same_v<decltype(binary_op(init, x)), T>,
827-
"Result type of binary_op must match scan accumulation type.");
828-
#else
829778
static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
830779
"Result type of binary_op must match scan accumulation type.");
831-
#endif
832-
833780
#ifdef __SYCL_DEVICE_ONLY__
834781
typename Group::linear_id_type local_linear_id =
835782
sycl::detail::get_local_linear_id(g);
@@ -884,17 +831,8 @@ std::enable_if_t<
884831
OutPtr>
885832
joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
886833
BinaryOperation binary_op) {
887-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
888-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
889-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
890-
? std::is_same_v<decltype(binary_op(init, *first)), bool>
891-
: std::is_same_v<decltype(binary_op(init, *first)), T>,
892-
"Result type of binary_op must match scan accumulation type.");
893-
#else
894834
static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
895835
"Result type of binary_op must match scan accumulation type.");
896-
#endif
897-
898836
#ifdef __SYCL_DEVICE_ONLY__
899837
ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
900838
ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
@@ -945,33 +883,9 @@ std::enable_if_t<
945883
OutPtr>
946884
joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
947885
BinaryOperation binary_op) {
948-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
949-
static_assert(
950-
(std::is_same_v<BinaryOperation,
951-
sycl::logical_or<std::remove_cv_t<
952-
std::remove_reference_t<decltype(*first)>>>> ||
953-
std::is_same_v<BinaryOperation,
954-
sycl::logical_and<std::remove_cv_t<
955-
std::remove_reference_t<decltype(*first)>>>>)
956-
? std::is_same_v<decltype(binary_op(
957-
std::remove_cv_t<
958-
std::remove_reference_t<decltype(*first)>>(),
959-
std::remove_cv_t<std::remove_reference_t<
960-
decltype(*first)>>())),
961-
bool>
962-
: std::is_same_v<
963-
decltype(binary_op(
964-
std::remove_cv_t<
965-
std::remove_reference_t<decltype(*first)>>(),
966-
std::remove_cv_t<
967-
std::remove_reference_t<decltype(*first)>>())),
968-
std::remove_cv_t<std::remove_reference_t<decltype(*first)>>>,
969-
"Result type of binary_op must match scan accumulation type.");
970-
#else
971886
static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
972887
typename detail::remove_pointer<OutPtr>::type>,
973888
"Result type of binary_op must match scan accumulation type.");
974-
#endif
975889
using T = typename detail::remove_pointer<OutPtr>::type;
976890
T init = detail::identity_for_ga_op<T, BinaryOperation>();
977891
return joint_exclusive_scan(g, first, last, result, init, binary_op);
@@ -989,19 +903,8 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
989903
detail::is_native_op<T, BinaryOperation>::value),
990904
T>
991905
inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
992-
993-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
994-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
995-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
996-
? std::is_same_v<decltype(binary_op(x, x)), bool>
997-
: std::is_same_v<decltype(binary_op(x, x)), T>,
998-
"Result type of binary_op must match scan accumulation type.");
999-
#else
1000-
1001906
static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
1002907
"Result type of binary_op must match scan accumulation type.");
1003-
#endif
1004-
1005908
#ifdef __SYCL_DEVICE_ONLY__
1006909
#if defined(__NVPTX__)
1007910
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<Group>) {
@@ -1069,18 +972,8 @@ std::enable_if_t<
1069972
std::is_convertible_v<V, T>),
1070973
T>
1071974
inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
1072-
1073-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
1074-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
1075-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
1076-
? std::is_same_v<decltype(binary_op(init, x)), bool>
1077-
: std::is_same_v<decltype(binary_op(init, x)), T>,
1078-
"Result type of binary_op must match scan accumulation type.");
1079-
#else
1080975
static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
1081976
"Result type of binary_op must match scan accumulation type.");
1082-
#endif
1083-
1084977
#ifdef __SYCL_DEVICE_ONLY__
1085978
T y = x;
1086979
if (sycl::detail::get_local_linear_id(g) == 0) {
@@ -1129,17 +1022,8 @@ std::enable_if_t<
11291022
OutPtr>
11301023
joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
11311024
BinaryOperation binary_op, T init) {
1132-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
1133-
static_assert((std::is_same_v<BinaryOperation, sycl::logical_or<T>> ||
1134-
std::is_same_v<BinaryOperation, sycl::logical_and<T>>)
1135-
? std::is_same_v<decltype(binary_op(init, *first)), bool>
1136-
: std::is_same_v<decltype(binary_op(init, *first)), T>,
1137-
"Result type of binary_op must match scan accumulation type.");
1138-
#else
11391025
static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
11401026
"Result type of binary_op must match scan accumulation type.");
1141-
#endif
1142-
11431027
#ifdef __SYCL_DEVICE_ONLY__
11441028
ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
11451029
ptrdiff_t stride = sycl::detail::get_local_linear_range(g);
@@ -1187,33 +1071,9 @@ std::enable_if_t<
11871071
OutPtr>
11881072
joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
11891073
BinaryOperation binary_op) {
1190-
#ifdef __INTEL_PREVIEW_BREAKING_CHANGES
1191-
static_assert(
1192-
(std::is_same_v<BinaryOperation,
1193-
sycl::logical_or<std::remove_cv_t<
1194-
std::remove_reference_t<decltype(*first)>>>> ||
1195-
std::is_same_v<BinaryOperation,
1196-
sycl::logical_and<std::remove_cv_t<
1197-
std::remove_reference_t<decltype(*first)>>>>)
1198-
? std::is_same_v<decltype(binary_op(
1199-
std::remove_cv_t<
1200-
std::remove_reference_t<decltype(*first)>>(),
1201-
std::remove_cv_t<std::remove_reference_t<
1202-
decltype(*first)>>())),
1203-
bool>
1204-
: std::is_same_v<
1205-
decltype(binary_op(
1206-
std::remove_cv_t<
1207-
std::remove_reference_t<decltype(*first)>>(),
1208-
std::remove_cv_t<
1209-
std::remove_reference_t<decltype(*first)>>())),
1210-
std::remove_cv_t<std::remove_reference_t<decltype(*first)>>>,
1211-
"Result type of binary_op must match scan accumulation type.");
1212-
#else
12131074
static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
12141075
typename detail::remove_pointer<OutPtr>::type>,
12151076
"Result type of binary_op must match scan accumulation type.");
1216-
#endif
12171077

12181078
using T = typename detail::remove_pointer<OutPtr>::type;
12191079
T init = detail::identity_for_ga_op<T, BinaryOperation>();

0 commit comments

Comments
 (0)