From 118830c1e53b317f9a4a8bd117bab6ff10aad928 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 12 May 2025 13:23:06 +0200 Subject: [PATCH 01/10] [SYCL] fix asserts after logical operation changes --- sycl/include/sycl/group_algorithm.hpp | 140 ++++++++++++++++++++++++++ 1 file changed, 140 insertions(+) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index 9547039d45b69..c475329340dc3 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -215,9 +215,19 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> reduce_over_group(Group g, T x, BinaryOperation binary_op) { + +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -291,9 +301,18 @@ std::enable_if_t< std::is_convertible_v), T> reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ return binary_op(init, reduce_over_group(g, T(x), binary_op)); #else @@ -341,9 +360,18 @@ std::enable_if_t< detail::is_native_op::value), T> joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ T partial = detail::identity_for_ga_op(); sycl::detail::for_each( @@ -679,8 +707,16 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -752,8 +788,16 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif T result; typename detail::get_scalar_binary_op::type scalar_binary_op{}; @@ -775,8 +819,17 @@ std::enable_if_t< std::is_convertible_v), T> exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ typename Group::linear_id_type local_linear_id = sycl::detail::get_local_linear_id(g); @@ -831,8 +884,17 @@ std::enable_if_t< OutPtr> joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ ptrdiff_t offset = sycl::detail::get_local_linear_id(g); ptrdiff_t stride = sycl::detail::get_local_linear_range(g); @@ -883,9 +945,33 @@ std::enable_if_t< OutPtr> joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert( + (std::is_same_v>>> || + std::is_same_v>>>) + ? std::is_same_v>(), + std::remove_cv_t>())), + bool> + : std::is_same_v< + decltype(binary_op( + std::remove_cv_t< + std::remove_reference_t>(), + std::remove_cv_t< + std::remove_reference_t>())), + std::remove_cv_t>>, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); +#endif using T = typename detail::remove_pointer::type; T init = detail::identity_for_ga_op(); return joint_exclusive_scan(g, first, last, result, init, binary_op); @@ -903,8 +989,19 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { + +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else + static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -972,8 +1069,18 @@ std::enable_if_t< std::is_convertible_v), T> inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) { + +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ T y = x; if (sycl::detail::get_local_linear_id(g) == 0) { @@ -1022,8 +1129,17 @@ std::enable_if_t< OutPtr> joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert((std::is_same_v> || + std::is_same_v>) + ? std::is_same_v + : std::is_same_v, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); +#endif + #ifdef __SYCL_DEVICE_ONLY__ ptrdiff_t offset = sycl::detail::get_local_linear_id(g); ptrdiff_t stride = sycl::detail::get_local_linear_range(g); @@ -1071,9 +1187,33 @@ std::enable_if_t< OutPtr> joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { +#ifdef __INTEL_PREVIEW_BREAKING_CHANGES + static_assert( + (std::is_same_v>>> || + std::is_same_v>>>) + ? std::is_same_v>(), + std::remove_cv_t>())), + bool> + : std::is_same_v< + decltype(binary_op( + std::remove_cv_t< + std::remove_reference_t>(), + std::remove_cv_t< + std::remove_reference_t>())), + std::remove_cv_t>>, + "Result type of binary_op must match scan accumulation type."); +#else static_assert(std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); +#endif using T = typename detail::remove_pointer::type; T init = detail::identity_for_ga_op(); From 3f34653efd74a0c78e1333946efc3003a34893eb Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 13 May 2025 11:59:59 +0200 Subject: [PATCH 02/10] [SYCL] use alias for readability --- sycl/include/sycl/group_algorithm.hpp | 50 ++++++++------------------- 1 file changed, 14 insertions(+), 36 deletions(-) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index c475329340dc3..808421e0540be 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -946,26 +946,15 @@ std::enable_if_t< joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { #ifdef __INTEL_PREVIEW_BREAKING_CHANGES + using binary_op_t = + std::remove_cv_t>; static_assert( - (std::is_same_v>>> || - std::is_same_v>>>) - ? std::is_same_v>(), - std::remove_cv_t>())), + (std::is_same_v> || + std::is_same_v>) + ? std::is_same_v - : std::is_same_v< - decltype(binary_op( - std::remove_cv_t< - std::remove_reference_t>(), - std::remove_cv_t< - std::remove_reference_t>())), - std::remove_cv_t>>, + : std::is_same_v, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v>; static_assert( - (std::is_same_v>>> || - std::is_same_v>>>) - ? std::is_same_v>(), - std::remove_cv_t>())), + (std::is_same_v> || + std::is_same_v>) + ? std::is_same_v - : std::is_same_v< - decltype(binary_op( - std::remove_cv_t< - std::remove_reference_t>(), - std::remove_cv_t< - std::remove_reference_t>())), - std::remove_cv_t>>, + : std::is_same_v, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v Date: Thu, 22 May 2025 15:06:00 +0200 Subject: [PATCH 03/10] [SYCL] use output type to check in assert --- sycl/include/sycl/group_algorithm.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index 808421e0540be..ad541fb17878d 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -953,8 +953,8 @@ joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, std::is_same_v>) ? std::is_same_v - : std::is_same_v, + : std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v>) ? std::is_same_v - : std::is_same_v, + : std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v Date: Mon, 13 Oct 2025 10:32:44 +0200 Subject: [PATCH 04/10] Revert "[SYCL] use output type to check in assert" This reverts commit f3b026aac5d859d7270c45fae479b62c04a18894. --- sycl/include/sycl/group_algorithm.hpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index ad541fb17878d..808421e0540be 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -953,8 +953,8 @@ joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, std::is_same_v>) ? std::is_same_v - : std::is_same_v::type>, + : std::is_same_v, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v>) ? std::is_same_v - : std::is_same_v::type>, + : std::is_same_v, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v Date: Mon, 13 Oct 2025 10:33:06 +0200 Subject: [PATCH 05/10] Revert "[SYCL] use alias for readability" This reverts commit 3f34653efd74a0c78e1333946efc3003a34893eb. --- sycl/include/sycl/group_algorithm.hpp | 50 +++++++++++++++++++-------- 1 file changed, 36 insertions(+), 14 deletions(-) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index 808421e0540be..c475329340dc3 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -946,15 +946,26 @@ std::enable_if_t< joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { #ifdef __INTEL_PREVIEW_BREAKING_CHANGES - using binary_op_t = - std::remove_cv_t>; static_assert( - (std::is_same_v> || - std::is_same_v>) - ? std::is_same_v>>> || + std::is_same_v>>>) + ? std::is_same_v>(), + std::remove_cv_t>())), bool> - : std::is_same_v, + : std::is_same_v< + decltype(binary_op( + std::remove_cv_t< + std::remove_reference_t>(), + std::remove_cv_t< + std::remove_reference_t>())), + std::remove_cv_t>>, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v>; static_assert( - (std::is_same_v> || - std::is_same_v>) - ? std::is_same_v>>> || + std::is_same_v>>>) + ? std::is_same_v>(), + std::remove_cv_t>())), bool> - : std::is_same_v, + : std::is_same_v< + decltype(binary_op( + std::remove_cv_t< + std::remove_reference_t>(), + std::remove_cv_t< + std::remove_reference_t>())), + std::remove_cv_t>>, "Result type of binary_op must match scan accumulation type."); #else static_assert(std::is_same_v Date: Mon, 13 Oct 2025 10:33:24 +0200 Subject: [PATCH 06/10] Revert "[SYCL] fix asserts after logical operation changes" This reverts commit 118830c1e53b317f9a4a8bd117bab6ff10aad928. --- sycl/include/sycl/group_algorithm.hpp | 140 -------------------------- 1 file changed, 140 deletions(-) diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp index c475329340dc3..9547039d45b69 100644 --- a/sycl/include/sycl/group_algorithm.hpp +++ b/sycl/include/sycl/group_algorithm.hpp @@ -215,19 +215,9 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> reduce_over_group(Group g, T x, BinaryOperation binary_op) { - -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -301,18 +291,9 @@ std::enable_if_t< std::is_convertible_v), T> reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ return binary_op(init, reduce_over_group(g, T(x), binary_op)); #else @@ -360,18 +341,9 @@ std::enable_if_t< detail::is_native_op::value), T> joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert( std::is_same_v, "Result type of binary_op must match reduction accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ T partial = detail::identity_for_ga_op(); sycl::detail::for_each( @@ -707,16 +679,8 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -788,16 +752,8 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif T result; typename detail::get_scalar_binary_op::type scalar_binary_op{}; @@ -819,17 +775,8 @@ std::enable_if_t< std::is_convertible_v), T> exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ typename Group::linear_id_type local_linear_id = sycl::detail::get_local_linear_id(g); @@ -884,17 +831,8 @@ std::enable_if_t< OutPtr> joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ ptrdiff_t offset = sycl::detail::get_local_linear_id(g); ptrdiff_t stride = sycl::detail::get_local_linear_range(g); @@ -945,33 +883,9 @@ std::enable_if_t< OutPtr> joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert( - (std::is_same_v>>> || - std::is_same_v>>>) - ? std::is_same_v>(), - std::remove_cv_t>())), - bool> - : std::is_same_v< - decltype(binary_op( - std::remove_cv_t< - std::remove_reference_t>(), - std::remove_cv_t< - std::remove_reference_t>())), - std::remove_cv_t>>, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); -#endif using T = typename detail::remove_pointer::type; T init = detail::identity_for_ga_op(); return joint_exclusive_scan(g, first, last, result, init, binary_op); @@ -989,19 +903,8 @@ std::enable_if_t<(is_group_v> && detail::is_native_op::value), T> inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) { - -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else - static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ #if defined(__NVPTX__) if constexpr (ext::oneapi::experimental::is_user_constructed_group_v) { @@ -1069,18 +972,8 @@ std::enable_if_t< std::is_convertible_v), T> inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) { - -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ T y = x; if (sycl::detail::get_local_linear_id(g) == 0) { @@ -1129,17 +1022,8 @@ std::enable_if_t< OutPtr> joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op, T init) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert((std::is_same_v> || - std::is_same_v>) - ? std::is_same_v - : std::is_same_v, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v, "Result type of binary_op must match scan accumulation type."); -#endif - #ifdef __SYCL_DEVICE_ONLY__ ptrdiff_t offset = sycl::detail::get_local_linear_id(g); ptrdiff_t stride = sycl::detail::get_local_linear_range(g); @@ -1187,33 +1071,9 @@ std::enable_if_t< OutPtr> joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, BinaryOperation binary_op) { -#ifdef __INTEL_PREVIEW_BREAKING_CHANGES - static_assert( - (std::is_same_v>>> || - std::is_same_v>>>) - ? std::is_same_v>(), - std::remove_cv_t>())), - bool> - : std::is_same_v< - decltype(binary_op( - std::remove_cv_t< - std::remove_reference_t>(), - std::remove_cv_t< - std::remove_reference_t>())), - std::remove_cv_t>>, - "Result type of binary_op must match scan accumulation type."); -#else static_assert(std::is_same_v::type>, "Result type of binary_op must match scan accumulation type."); -#endif using T = typename detail::remove_pointer::type; T init = detail::identity_for_ga_op(); From aaaf763c255afc5ed634ff457fc68a8e1a089a0e Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 13 Oct 2025 11:05:07 +0200 Subject: [PATCH 07/10] [SYCL] do not use logical_or and logical_and in group algorithm e2e tests --- .../GroupAlgorithm/exclusive_scan_sycl2020.cpp | 12 ------------ .../GroupAlgorithm/inclusive_scan_sycl2020.cpp | 17 ----------------- .../test-e2e/GroupAlgorithm/reduce_sycl2020.cpp | 12 ------------ 3 files changed, 41 deletions(-) diff --git a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 8d60cea9a377d..c1142703b8368 100644 --- a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -1,9 +1,6 @@ // RUN: %{build} -Wno-error=deprecated-declarations -fsycl-device-code-split=per_kernel -I . -o %t.out // RUN: %{run} %t.out -// XFAIL: preview-mode -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18390 - #include "../helpers.hpp" #include "support.h" #include @@ -173,15 +170,6 @@ int main() { test(input, sycl::bit_xor(), 0); test(input_small, sycl::bit_and(), ~0); - test(input, sycl::logical_or(), 0); - test(input, sycl::logical_and(), 1); - - std::array bool_input = {}; - test(bool_input, sycl::logical_or(), false); - test(bool_input, sycl::logical_or<>(), false); - test(bool_input, sycl::logical_and(), true); - test(bool_input, sycl::logical_and<>(), true); - std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); test(int2_input, sycl::plus(), {0, 0}); diff --git a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 73fd739b23a6a..dbba805f12b2f 100644 --- a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -1,9 +1,6 @@ // RUN: %{build} -Wno-error=deprecated-declarations -fsycl-device-code-split=per_kernel -I . -o %t.out // RUN: %{run} %t.out -// XFAIL: preview-mode -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18390 - #include "../helpers.hpp" #include "support.h" #include @@ -171,20 +168,6 @@ int main() { test(input, sycl::bit_xor(), 0); test(input_small, sycl::bit_and(), ~0); - test(input, sycl::logical_or(), 0); - test(input, sycl::logical_and(), 1); - - std::array bool_input = {}; - test(bool_input, sycl::logical_or(), false); - test(bool_input, sycl::logical_or<>(), false); - test(bool_input, sycl::logical_and(), true); - test(bool_input, sycl::logical_and<>(), true); - - test(bool_input, sycl::logical_or(), false); - test(bool_input, sycl::logical_or<>(), false); - test(bool_input, sycl::logical_and(), true); - test(bool_input, sycl::logical_and<>(), true); - std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); test(int2_input, sycl::plus(), {0, 0}); diff --git a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp index ec12db383c8c1..cc2f8a310bbee 100644 --- a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp @@ -1,9 +1,6 @@ // RUN: %{build} -fsycl-device-code-split=per_kernel -I . -o %t.out // RUN: %{run} %t.out -// XFAIL: preview-mode -// XFAIL-TRACKER: https://github.com/intel/llvm/issues/18390 - #include "support.h" #include @@ -100,15 +97,6 @@ int main() { test(input, sycl::bit_xor(), 0); test(input, sycl::bit_and(), ~0); - test(input, sycl::logical_or(), 0); - test(input, sycl::logical_and(), 1); - - std::array bool_input = {}; - test(bool_input, sycl::logical_or(), false); - test(bool_input, sycl::logical_or<>(), false); - test(bool_input, sycl::logical_and(), true); - test(bool_input, sycl::logical_and<>(), true); - std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); test(int2_input, sycl::plus(), {0, 0}); From 20c6e47d0c4826ff69dbfcb41e89adeca927ae5b Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 16 Oct 2025 16:19:00 +0200 Subject: [PATCH 08/10] [SYCL] implement tests to check static assert of group algorithms --- .../exclusive_scan_sycl2020.cpp | 4 + .../inclusive_scan_sycl2020.cpp | 5 +- .../logical_or_and_group_algorithms.cpp | 174 ++++++++++++++++++ .../GroupAlgorithm/reduce_sycl2020.cpp | 4 + 4 files changed, 186 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp diff --git a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp index c1142703b8368..cff6376edb723 100644 --- a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp @@ -170,6 +170,10 @@ int main() { test(input, sycl::bit_xor(), 0); test(input_small, sycl::bit_and(), ~0); + std::array bool_input = {}; + test(bool_input, sycl::logical_or(), false); + test(bool_input, sycl::logical_and(), true); + std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); test(int2_input, sycl::plus(), {0, 0}); diff --git a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp index dbba805f12b2f..7e0d6aeb509f4 100644 --- a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp @@ -166,7 +166,10 @@ int main() { test(input_small, sycl::multiplies(), 1); test(input, sycl::bit_or(), 0); test(input, sycl::bit_xor(), 0); - test(input_small, sycl::bit_and(), ~0); + + std::array bool_input = {}; + test(bool_input, sycl::logical_or(), false); + test(bool_input, sycl::logical_and(), true); std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); diff --git a/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp b/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp new file mode 100644 index 0000000000000..b80bda5a7606d --- /dev/null +++ b/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp @@ -0,0 +1,174 @@ +// RUN: %clangxx -fsycl -Xclang -verify=expected -fpreview-breaking-changes -fsyntax-only -ferror-limit=0 %s + +// expected-error@sycl/group_algorithm.hpp:* 16 {{Result type of binary_op must match scan accumulation type}} +// expected-error@sycl/group_algorithm.hpp:* 6 {{Result type of binary_op must match reduction accumulation type}} + +#include +#include +#include +#include +#include + +using namespace sycl; + +void TestExclusiveScanOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, sycl::logical_and>' requested here}} + exclusive_scan_over_group(g, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, int, sycl::logical_and>' requested here}} + exclusive_scan_over_group(g, 0, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, sycl::logical_or>' requested here}} + exclusive_scan_over_group(g, 0, sycl::logical_or{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, int, sycl::logical_or>' requested here}} + exclusive_scan_over_group(g, 0, 0, sycl::logical_or{}); + }); + }); +} + +void TestJointExclusiveScan(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and>' requested here}} + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or>' requested here}} + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}); + + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, int, sycl::logical_and>' requested here}} + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, + sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, int, sycl::logical_or>' requested here}} + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, + sycl::logical_or{}); + }); + }).wait(); +} + +void TestInclusiveScanOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_and>' requested here}} + inclusive_scan_over_group(g, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_and, int>' requested here}} + inclusive_scan_over_group(g, 0, sycl::logical_and{}, 0); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_or>' requested here}} + inclusive_scan_over_group(g, 0, sycl::logical_or{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_or, int>' requested here}} + inclusive_scan_over_group(g, 0, sycl::logical_or{}, 0); + }); + }); +} + +void TestJointInclusiveScan(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and>' requested here}} + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or>' requested here}} + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}); + + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and, int>' requested here}} + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}, 0); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or, int>' requested here}} + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}, 0); + }); + }).wait(); +} + + +void TestReduceOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, sycl::logical_and>' requested here}} + reduce_over_group(g, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, int, sycl::logical_and>' requested here}} + reduce_over_group(g, 0, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, sycl::logical_or>' requested here}} + reduce_over_group(g, 0, sycl::logical_or{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, int, sycl::logical_or>' requested here}} + reduce_over_group(g, 0, 0, sycl::logical_or{}); + }); + }); +} + +void TestJointReduce(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_reduce, sycl::multi_ptr, int, sycl::logical_and>' requested here}} + joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_and{}); + // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_reduce, sycl::multi_ptr, int, sycl::logical_or>' requested here}} + joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_or{}); + }); + }).wait(); +} + + +int main() { + sycl::queue q; + TestExclusiveScanOverGroup(q); + TestJointExclusiveScan(q); + TestInclusiveScanOverGroup(q); + TestJointInclusiveScan(q); + TestReduceOverGroup(q); + TestJointReduce(q); + return 0; +} diff --git a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp index cc2f8a310bbee..9bc40d9a2c909 100644 --- a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp +++ b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp @@ -97,6 +97,10 @@ int main() { test(input, sycl::bit_xor(), 0); test(input, sycl::bit_and(), ~0); + std::array bool_input = {}; + test(bool_input, sycl::logical_or(), false); + test(bool_input, sycl::logical_and(), true); + std::array int2_input = {}; std::iota(int2_input.begin(), int2_input.end(), 0); test(int2_input, sycl::plus(), {0, 0}); From eabe3b75355875a6ec3cc182350976340d1bc985 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 16 Oct 2025 16:26:28 +0200 Subject: [PATCH 09/10] [SYCL] fix formatting in tests --- .../GroupAlgorithm/logical_or_and_group_algorithms.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp b/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp index b80bda5a7606d..afbdf0127dbc1 100644 --- a/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp +++ b/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp @@ -3,11 +3,11 @@ // expected-error@sycl/group_algorithm.hpp:* 16 {{Result type of binary_op must match scan accumulation type}} // expected-error@sycl/group_algorithm.hpp:* 6 {{Result type of binary_op must match reduction accumulation type}} -#include +#include +#include #include #include -#include -#include +#include using namespace sycl; @@ -116,7 +116,6 @@ void TestJointInclusiveScan(sycl::queue &q) { }).wait(); } - void TestReduceOverGroup(sycl::queue &q) { q.submit([&](handler &cgh) { cgh.parallel_for( @@ -161,7 +160,6 @@ void TestJointReduce(sycl::queue &q) { }).wait(); } - int main() { sycl::queue q; TestExclusiveScanOverGroup(q); From fc07225fed78bd9439b3006a0c3221f05c8cf98d Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 20 Oct 2025 10:42:40 +0200 Subject: [PATCH 10/10] [SYCL] move static assert checks from tests-e2e folder to test --- .../logical_or_and_group_algorithms.cpp | 172 ------------------ .../logical_or_and_group_algorithms.cpp | 148 +++++++++++++++ 2 files changed, 148 insertions(+), 172 deletions(-) delete mode 100644 sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp create mode 100644 sycl/test/group_algorithms/logical_or_and_group_algorithms.cpp diff --git a/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp b/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp deleted file mode 100644 index afbdf0127dbc1..0000000000000 --- a/sycl/test-e2e/GroupAlgorithm/logical_or_and_group_algorithms.cpp +++ /dev/null @@ -1,172 +0,0 @@ -// RUN: %clangxx -fsycl -Xclang -verify=expected -fpreview-breaking-changes -fsyntax-only -ferror-limit=0 %s - -// expected-error@sycl/group_algorithm.hpp:* 16 {{Result type of binary_op must match scan accumulation type}} -// expected-error@sycl/group_algorithm.hpp:* 6 {{Result type of binary_op must match reduction accumulation type}} - -#include -#include -#include -#include -#include - -using namespace sycl; - -void TestExclusiveScanOverGroup(sycl::queue &q) { - q.submit([&](handler &cgh) { - cgh.parallel_for( - nd_range<1>(1, 1), [=](nd_item<1> it) { - group<1> g = it.get_group(); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, sycl::logical_and>' requested here}} - exclusive_scan_over_group(g, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, int, sycl::logical_and>' requested here}} - exclusive_scan_over_group(g, 0, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, sycl::logical_or>' requested here}} - exclusive_scan_over_group(g, 0, sycl::logical_or{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::exclusive_scan_over_group, int, int, sycl::logical_or>' requested here}} - exclusive_scan_over_group(g, 0, 0, sycl::logical_or{}); - }); - }); -} - -void TestJointExclusiveScan(sycl::queue &q) { - constexpr size_t N = 8; - int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; - int output[N] = {}; - - sycl::buffer inBuf(input, sycl::range<1>(N)); - sycl::buffer outBuf(output, sycl::range<1>(N)); - - q.submit([&](sycl::handler &cgh) { - auto in = inBuf.get_access(cgh); - auto out = outBuf.get_access(cgh); - - cgh.parallel_for( - sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), - [=](sycl::nd_item<1> it) { - auto g = it.get_group(); - auto inPtr = in.get_multi_ptr(); - auto outPtr = out.get_multi_ptr(); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and>' requested here}} - joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or>' requested here}} - joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_or{}); - - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, int, sycl::logical_and>' requested here}} - joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, - sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_exclusive_scan, sycl::multi_ptr, sycl::multi_ptr, int, sycl::logical_or>' requested here}} - joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, - sycl::logical_or{}); - }); - }).wait(); -} - -void TestInclusiveScanOverGroup(sycl::queue &q) { - q.submit([&](handler &cgh) { - cgh.parallel_for( - nd_range<1>(1, 1), [=](nd_item<1> it) { - group<1> g = it.get_group(); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_and>' requested here}} - inclusive_scan_over_group(g, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_and, int>' requested here}} - inclusive_scan_over_group(g, 0, sycl::logical_and{}, 0); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_or>' requested here}} - inclusive_scan_over_group(g, 0, sycl::logical_or{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::inclusive_scan_over_group, int, sycl::logical_or, int>' requested here}} - inclusive_scan_over_group(g, 0, sycl::logical_or{}, 0); - }); - }); -} - -void TestJointInclusiveScan(sycl::queue &q) { - constexpr size_t N = 8; - int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; - int output[N] = {}; - - sycl::buffer inBuf(input, sycl::range<1>(N)); - sycl::buffer outBuf(output, sycl::range<1>(N)); - - q.submit([&](sycl::handler &cgh) { - auto in = inBuf.get_access(cgh); - auto out = outBuf.get_access(cgh); - - cgh.parallel_for( - sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), - [=](sycl::nd_item<1> it) { - auto g = it.get_group(); - auto inPtr = in.get_multi_ptr(); - auto outPtr = out.get_multi_ptr(); - - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and>' requested here}} - joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or>' requested here}} - joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_or{}); - - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_and, int>' requested here}} - joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_and{}, 0); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_inclusive_scan, sycl::multi_ptr, sycl::multi_ptr, sycl::logical_or, int>' requested here}} - joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, - sycl::logical_or{}, 0); - }); - }).wait(); -} - -void TestReduceOverGroup(sycl::queue &q) { - q.submit([&](handler &cgh) { - cgh.parallel_for( - nd_range<1>(1, 1), [=](nd_item<1> it) { - group<1> g = it.get_group(); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, sycl::logical_and>' requested here}} - reduce_over_group(g, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, int, sycl::logical_and>' requested here}} - reduce_over_group(g, 0, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, sycl::logical_or>' requested here}} - reduce_over_group(g, 0, sycl::logical_or{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::reduce_over_group, int, int, sycl::logical_or>' requested here}} - reduce_over_group(g, 0, 0, sycl::logical_or{}); - }); - }); -} - -void TestJointReduce(sycl::queue &q) { - constexpr size_t N = 8; - int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; - int output[N] = {}; - - sycl::buffer inBuf(input, sycl::range<1>(N)); - sycl::buffer outBuf(output, sycl::range<1>(N)); - - q.submit([&](sycl::handler &cgh) { - auto in = inBuf.get_access(cgh); - auto out = outBuf.get_access(cgh); - - cgh.parallel_for( - sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), - [=](sycl::nd_item<1> it) { - auto g = it.get_group(); - auto inPtr = in.get_multi_ptr(); - auto outPtr = out.get_multi_ptr(); - - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_reduce, sycl::multi_ptr, int, sycl::logical_and>' requested here}} - joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_and{}); - // expected-note@+1 {{in instantiation of function template specialization 'sycl::joint_reduce, sycl::multi_ptr, int, sycl::logical_or>' requested here}} - joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_or{}); - }); - }).wait(); -} - -int main() { - sycl::queue q; - TestExclusiveScanOverGroup(q); - TestJointExclusiveScan(q); - TestInclusiveScanOverGroup(q); - TestJointInclusiveScan(q); - TestReduceOverGroup(q); - TestJointReduce(q); - return 0; -} diff --git a/sycl/test/group_algorithms/logical_or_and_group_algorithms.cpp b/sycl/test/group_algorithms/logical_or_and_group_algorithms.cpp new file mode 100644 index 0000000000000..05f9bc996c21b --- /dev/null +++ b/sycl/test/group_algorithms/logical_or_and_group_algorithms.cpp @@ -0,0 +1,148 @@ +// RUN: %clangxx -fsycl -Xclang -verify=expected -Xclang -verify-ignore-unexpected=note -fpreview-breaking-changes -fsyntax-only -ferror-limit=0 %s + +// expected-error@sycl/group_algorithm.hpp:* 16 {{Result type of binary_op must match scan accumulation type}} +// expected-error@sycl/group_algorithm.hpp:* 6 {{Result type of binary_op must match reduction accumulation type}} + +#include +#include +#include +#include +#include + +using namespace sycl; + +void TestExclusiveScanOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + exclusive_scan_over_group(g, 0, sycl::logical_and{}); + exclusive_scan_over_group(g, 0, 0, sycl::logical_and{}); + exclusive_scan_over_group(g, 0, sycl::logical_or{}); + exclusive_scan_over_group(g, 0, 0, sycl::logical_or{}); + }); + }); +} + +void TestJointExclusiveScan(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}); + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}); + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, + sycl::logical_and{}); + joint_exclusive_scan(g, inPtr, inPtr + N, outPtr, 0, + sycl::logical_or{}); + }); + }).wait(); +} + +void TestInclusiveScanOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + inclusive_scan_over_group(g, 0, sycl::logical_and{}); + inclusive_scan_over_group(g, 0, sycl::logical_and{}, 0); + inclusive_scan_over_group(g, 0, sycl::logical_or{}); + inclusive_scan_over_group(g, 0, sycl::logical_or{}, 0); + }); + }); +} + +void TestJointInclusiveScan(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}); + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}); + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_and{}, 0); + joint_inclusive_scan(g, inPtr, inPtr + N, outPtr, + sycl::logical_or{}, 0); + }); + }).wait(); +} + +void TestReduceOverGroup(sycl::queue &q) { + q.submit([&](handler &cgh) { + cgh.parallel_for( + nd_range<1>(1, 1), [=](nd_item<1> it) { + group<1> g = it.get_group(); + reduce_over_group(g, 0, sycl::logical_and{}); + reduce_over_group(g, 0, 0, sycl::logical_and{}); + reduce_over_group(g, 0, sycl::logical_or{}); + reduce_over_group(g, 0, 0, sycl::logical_or{}); + }); + }); +} + +void TestJointReduce(sycl::queue &q) { + constexpr size_t N = 8; + int input[N] = {1, 2, 3, 4, 5, 6, 7, 8}; + int output[N] = {}; + + sycl::buffer inBuf(input, sycl::range<1>(N)); + sycl::buffer outBuf(output, sycl::range<1>(N)); + + q.submit([&](sycl::handler &cgh) { + auto in = inBuf.get_access(cgh); + auto out = outBuf.get_access(cgh); + + cgh.parallel_for( + sycl::nd_range<1>(sycl::range<1>(N), sycl::range<1>(N)), + [=](sycl::nd_item<1> it) { + auto g = it.get_group(); + auto inPtr = in.get_multi_ptr(); + auto outPtr = out.get_multi_ptr(); + + joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_and{}); + joint_reduce(g, inPtr, inPtr + N, 0, sycl::logical_or{}); + }); + }).wait(); +} + +int main() { + sycl::queue q; + TestExclusiveScanOverGroup(q); + TestJointExclusiveScan(q); + TestInclusiveScanOverGroup(q); + TestJointInclusiveScan(q); + TestReduceOverGroup(q); + TestJointReduce(q); + return 0; +}