diff --git a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp index 8d60cea9a377d..cff6376edb723 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,14 +170,9 @@ 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); diff --git a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp index 73fd739b23a6a..7e0d6aeb509f4 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 @@ -169,21 +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); - - 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); diff --git a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp index ec12db383c8c1..9bc40d9a2c909 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,14 +97,9 @@ 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); 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; +}