[Testing] Split up scan tests

test/parallel_api/numeric/numeric.ops/exclusive_scan.pass.cpp

@@ -23,9 +23,71 @@
 #include "support/scan_serial_impl.h"
 
 #include <iostream>
+#include <cstdint>
 #include <vector>
 
-#include "support/scan_serial_impl.h"
+using namespace TestUtils;
+
+template <typename In, typename Init, typename Out>
+struct test_exclusive_scan_with_plus
+{
+    template <typename Policy, typename Iterator1, typename Iterator2, typename Iterator3, typename Size, typename T>
+    std::enable_if_t<!TestUtils::is_reverse_v<Iterator1> || std::is_same_v<Iterator1, Iterator2>>
+    operator()(Policy&& exec, Iterator1 in_first, Iterator1 in_last, Iterator2 out_first, Iterator2 out_last,
+               Iterator3 expected_first, Iterator3 /* expected_last */, Size n, T init, T trash)
+    {
+        using namespace std;
+
+        exclusive_scan_serial(in_first, in_last, expected_first, init);
+        auto orr = exclusive_scan(std::forward<Policy>(exec), in_first, in_last, out_first, init);
+        EXPECT_TRUE(out_last == orr, "exclusive_scan returned wrong iterator");
+        EXPECT_EQ_N(expected_first, out_first, n, "wrong result from exclusive_scan");
+        std::fill_n(out_first, n, trash);
+    }
+    // exclusive_scan with reverse_iterator between different iterator types results in a compilation error even if
+    // the call should be valid. Please see: https://github.com/uxlfoundation/oneDPL/issues/2296
+    template <typename Policy, typename Iterator1, typename Iterator2, typename Iterator3, typename Size, typename T>
+    std::enable_if_t<TestUtils::is_reverse_v<Iterator1> && !std::is_same_v<Iterator1, Iterator2>>
+    operator()(Policy&& /*exec*/, Iterator1 /*in_first*/, Iterator1 /*in_last*/, Iterator2 /*out_first*/,
+               Iterator2 /*out_last*/, Iterator3 /*expected_first*/, Iterator3 /*expected_last*/, Size /*n*/,
+               T /*init*/, T /*trash*/)
+    {
+    }
+};
+
+template <typename In, typename Init, typename Out, typename Convert>
+void
+test_with_plus(Init init, Out trash, Convert convert)
+{
+    for (size_t n = 0; n <= 100000; n = n <= 16 ? n + 1 : size_t(3.1415 * n))
+    {
+        Sequence<In> in(n, convert);
+        Sequence<Out> expected(n);
+        Sequence<Out> out(n, [&](std::int32_t) { return trash; });
+
+        invoke_on_all_policies<2>()(test_exclusive_scan_with_plus<In, Init, Out>(), in.begin(), in.end(), out.begin(),
+                                    out.end(), expected.begin(), expected.end(), in.size(), init, trash);
+        invoke_on_all_policies<3>()(test_exclusive_scan_with_plus<In, Init, Out>(), in.cbegin(), in.cend(), out.begin(),
+                                    out.end(), expected.begin(), expected.end(), in.size(), init, trash);
+    }
+
+#if TEST_DPCPP_BACKEND_PRESENT && !ONEDPL_FPGA_DEVICE
+    // testing of large number of items may take too much time in debug mode
+    unsigned long n =
+#if PSTL_USE_DEBUG
+        1000000;
+#else
+        100000000;
+#endif
+
+    Sequence<In> in(n, convert);
+    Sequence<Out> expected(n);
+    Sequence<Out> out(n, [&](std::int32_t) { return trash; });
+    invoke_on_all_hetero_policies<5>()(test_exclusive_scan_with_plus<In, Init, Out>(), in.begin(), in.end(),
+                                       out.begin(), out.end(), expected.begin(), expected.end(), in.size(), init,
+                                       trash);
+#endif // TEST_DPCPP_BACKEND_PRESENT && !ONEDPL_FPGA_DEVICE
+}
 
 #if TEST_DPCPP_BACKEND_PRESENT
 
@@ -68,7 +130,7 @@ template <sycl::usm::alloc alloc_type, typename Policy>
 void
 test_with_usm(Policy&& exec)
 {
-    for (::std::size_t n = 0; n <= TestUtils::max_n; n = n <= 16 ? n + 1 : size_t(3.1415 * n))
+    for (std::size_t n = 0; n <= TestUtils::max_n; n = n <= 16 ? n + 1 : size_t(3.1415 * n))
     {
         test_with_usm<alloc_type>(CLONE_TEST_POLICY(exec), n);
     }
@@ -81,11 +143,11 @@ test_diff_iterators(Policy&& exec)
     constexpr std::size_t N = 6;
 
     sycl::queue q = exec.queue();
-    
+
     // Allocate USM shared memory for input (bool type) and output (int type)
     bool* input = sycl::malloc_shared<bool>(N, q);
     int* result = sycl::malloc_shared<int>(N, q);
-    
+
     // Initialize input data
     input[0] = true;
     input[1] = false;
@@ -131,7 +193,8 @@ test_diff_iterators(Policy&& exec)
 }
 
 template <typename Policy>
-void test_impl(Policy&& exec)
+void
+test_usm_impl(Policy&& exec)
 {
     // Run tests for USM shared/device memory
     test_with_usm<sycl::usm::alloc::shared>(CLONE_TEST_POLICY(exec));
@@ -144,15 +207,25 @@ void test_impl(Policy&& exec)
 int
 main()
 {
-#if TEST_DPCPP_BACKEND_PRESENT
+    // Since the implicit "+" forms of the scan delegate to the generic forms,
+    // there's little point in using a highly restricted type, so just use double.
+    test_with_plus<float64_t, float64_t, float64_t>(
+        0.0, -666.0, [](std::uint32_t k) { return float64_t((k % 991 + 1) ^ (k % 997 + 2)); });
+    test_with_plus<std::int32_t, std::int32_t, std::int32_t>(
+        0.0, -666.0, [](std::uint32_t k) { return std::int32_t((k % 991 + 1) ^ (k % 997 + 2)); });
 
+    // When testing from bool to uint32_t, we must give a uint32_t init type to scan over integers
+    test_with_plus<bool, std::uint32_t, std::uint32_t>(0, 123456,
+                                                       [](std::uint32_t k) { return std::uint32_t{k % 2 == 0}; });
+
+#if TEST_DPCPP_BACKEND_PRESENT
     auto policy = TestUtils::get_dpcpp_test_policy();
-    test_impl(policy);
+    test_usm_impl(policy);
 
-#if TEST_CHECK_COMPILATION_WITH_DIFF_POLICY_VAL_CATEGORY
-    TestUtils::check_compilation(policy, [](auto&& policy) { test_impl(std::forward<decltype(policy)>(policy)); });
-#endif
+#    if TEST_CHECK_COMPILATION_WITH_DIFF_POLICY_VAL_CATEGORY
+    TestUtils::check_compilation(policy, [](auto&& policy) { test_usm_impl(std::forward<decltype(policy)>(policy)); });
+#    endif
 #endif // TEST_DPCPP_BACKEND_PRESENT
 
-    return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT);
+    return done();
 }

test/parallel_api/numeric/numeric.ops/exclusive_scan_matrix.pass.cpp

@@ -0,0 +1,62 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Copyright (C) UXL Foundation Contributors
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "support/test_config.h"
+
+#include _PSTL_TEST_HEADER(execution)
+#include _PSTL_TEST_HEADER(algorithm)
+#include _PSTL_TEST_HEADER(numeric)
+
+#include "support/utils.h"
+#include "support/scan_serial_impl.h"
+#include "support/utils_scan.h"
+
+#include <cstdint>
+
+using namespace TestUtils;
+
+template <typename In, typename Out, typename BinaryOp>
+void
+test_matrix(Out init, BinaryOp binary_op, Out trash)
+{
+    for (size_t n = 0; n <= 100000; n = n <= 16 ? n + 1 : size_t(3.1415 * n))
+    {
+        Sequence<In> in(n, [](size_t k) { return In(k, k + 1); });
+
+        Sequence<Out> out(n, [&](size_t) { return trash; });
+        Sequence<Out> expected(n, [&](size_t) { return trash; });
+
+        auto __scan_invoker = [&](Sequence<Out>& out) {
+#if !TEST_GCC10_EXCLUSIVE_SCAN_BROKEN
+            invoke_on_all_policies<8>()(test_exclusive_scan_with_binary_op<In>(), in.begin(), in.end(), out.begin(),
+                                        out.end(), expected.begin(), expected.end(), in.size(), init, binary_op, trash);
+            invoke_on_all_policies<9>()(test_exclusive_scan_with_binary_op<In>(), in.cbegin(), in.cend(), out.begin(),
+                                        out.end(), expected.begin(), expected.end(), in.size(), init, binary_op, trash);
+#endif
+        };
+
+        //perform regular a scan algorithm
+        __scan_invoker(out);
+
+        //perform an in-place scan algorithm
+        __scan_invoker(in);
+    }
+}
+
+int
+main()
+{
+#if !_PSTL_ICC_19_TEST_SIMD_UDS_WINDOWS_RELEASE_BROKEN
+    // Test with highly restricted type and associative but not commutative operation
+    test_matrix<Matrix2x2<std::int32_t>, Matrix2x2<std::int32_t>>(
+        Matrix2x2<std::int32_t>(), multiply_matrix<std::int32_t>(), Matrix2x2<std::int32_t>(-666, 666));
+#endif
+
+    return done();
+}

test/parallel_api/numeric/numeric.ops/exclusive_scan_multiplies.pass.cpp

@@ -0,0 +1,60 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Copyright (C) UXL Foundation Contributors
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "support/test_config.h"
+
+#include _PSTL_TEST_HEADER(execution)
+#include _PSTL_TEST_HEADER(algorithm)
+#include _PSTL_TEST_HEADER(numeric)
+
+#include "support/utils.h"
+#include "support/scan_serial_impl.h"
+#include "support/utils_scan.h"
+
+#include <random>
+#include <algorithm>
+#include <cstdint>
+#include <vector>
+
+using namespace TestUtils;
+
+template <typename T>
+void
+test_with_multiplies()
+{
+#if TEST_DPCPP_BACKEND_PRESENT
+    T trash = 666;
+    T init = 1;
+    const std::size_t custom_item_count = 10;
+
+    for (size_t n = custom_item_count; n <= 100000; n = n <= 16 ? n + 1 : size_t(3.1415 * n))
+    {
+        Sequence<T> out(n, [&](size_t) { return trash; });
+        Sequence<T> expected(n, [&](size_t) { return trash; });
+
+        Sequence<T> in(n, [](size_t /*index*/) { return 1; });
+        std::size_t counter = 0;
+        std::generate_n(in.begin(), custom_item_count, [&counter]() { return (counter++) % 3 + 2; });
+        std::default_random_engine gen{42};
+        std::shuffle(in.begin(), in.end(), gen);
+
+        invoke_on_all_hetero_policies<21>()(test_exclusive_scan_with_binary_op<T>(), in.begin(), in.end(), out.begin(),
+                                            out.end(), expected.begin(), expected.end(), in.size(), init,
+                                            std::multiplies{}, trash);
+    }
+#endif // TEST_DPCPP_BACKEND_PRESENT
+}
+
+int
+main()
+{
+    test_with_multiplies<std::uint64_t>();
+
+    return done(TEST_DPCPP_BACKEND_PRESENT);
+}

test/parallel_api/numeric/numeric.ops/exclusive_scan_usm.pass.cpp

@@ -0,0 +1,104 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Copyright (C) UXL Foundation Contributors
+//
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "support/test_config.h"
+
+#include "oneapi/dpl/execution"
+#include "oneapi/dpl/numeric"
+#include "oneapi/dpl/iterator"
+
+#include "support/utils.h"
+#include "support/scan_serial_impl.h"
+#include "support/utils_scan.h"
+
+using namespace TestUtils;
+
+#if TEST_DPCPP_BACKEND_PRESENT
+static const char kMsgExclusiveScanNormal [] = "Wrong effect from exclusive scan (non-inplace)";
+static const char kMsgExclusiveScanInplace[] = "Wrong effect from exclusive scan (inplace)";
+
+template <int InitValue>
+struct TestingAlgoritmExclusiveScan
+{
+    template <typename... TArgs>
+    void
+    call_onedpl(TArgs&&... args)
+    {
+        oneapi::dpl::exclusive_scan(std::forward<TArgs>(args)..., InitValue);
+    }
+
+    template <typename... TArgs>
+    void
+    call_serial(TArgs&&... args)
+    {
+        exclusive_scan_serial(std::forward<TArgs>(args)..., InitValue);
+    }
+
+    const char*
+    getMsg(bool bInplace) const
+    {
+        return bInplace ? kMsgExclusiveScanInplace : kMsgExclusiveScanNormal;
+    }
+};
+
+template <int InitValue, typename BinaryOp>
+struct TestingAlgoritmExclusiveScanExt
+{
+    template <typename... TArgs>
+    void
+    call_onedpl(TArgs&&... args)
+    {
+        oneapi::dpl::exclusive_scan(std::forward<TArgs>(args)..., InitValue, BinaryOp());
+    }
+
+    template <typename... TArgs>
+    void
+    call_serial(TArgs&&... args)
+    {
+        exclusive_scan_serial(std::forward<TArgs>(args)..., InitValue, BinaryOp());
+    }
+
+    const char*
+    getMsg(bool bInplace) const
+    {
+        return bInplace ? kMsgExclusiveScanInplace : kMsgExclusiveScanNormal;
+    }
+};
+
+template <sycl::usm::alloc alloc_type, typename ValueType, typename BinaryOperation>
+void
+run_test()
+{
+    // Non inplace
+    test2buffers<alloc_type, test_scan_non_inplace<ValueType, TestingAlgoritmExclusiveScan<2>>>();
+    test2buffers<alloc_type, test_scan_non_inplace<ValueType, TestingAlgoritmExclusiveScanExt<2, BinaryOperation>>>();
+
+    // Inplace
+    test1buffer<alloc_type, test_scan_inplace<ValueType, TestingAlgoritmExclusiveScan<2>>>();
+    test1buffer<alloc_type, test_scan_inplace<ValueType, TestingAlgoritmExclusiveScanExt<2, BinaryOperation>>>();
+}
+#endif // TEST_DPCPP_BACKEND_PRESENT
+
+int
+main()
+{
+#if TEST_DPCPP_BACKEND_PRESENT
+
+    using ValueType = int;
+    using BinaryOperation = std::plus<ValueType>;
+
+    // Run tests for USM shared memory
+    run_test<sycl::usm::alloc::shared, ValueType, BinaryOperation>();
+    // Run tests for USM device memory
+    run_test<sycl::usm::alloc::device, ValueType, BinaryOperation>();
+
+#endif // TEST_DPCPP_BACKEND_PRESENT
+
+    return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT);
+}