diff --git a/test/parallel_api/iterator/input_data_sweep.h b/test/parallel_api/iterator/input_data_sweep.h new file mode 100644 index 00000000000..fa2c028d48e --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep.h @@ -0,0 +1,224 @@ +// -*- C++ -*- +//===-- input_data_sweep.h ------------------------------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#ifndef _INPUT_DATA_SWEEP_H +#define _INPUT_DATA_SWEEP_H + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "support/utils_invoke.h" + +#if TEST_DPCPP_BACKEND_PRESENT +template +void +wrap_recurse(Policy&& exec, InputIterator1 first, InputIterator1 last, InputIterator2 copy_from_first, + OutputIterator copy_to_first, OriginalIterator1 orig_first, OriginalIterator2 orig_out_first, + ExpectedIterator expected_first, T trash, const std::string& input_descr) +{ + auto exec1 = TestUtils::create_new_policy_idx<0>(exec); + auto exec2 = TestUtils::create_new_policy_idx<1>(exec); + auto exec3 = TestUtils::create_new_policy_idx<2>(exec); + auto exec4 = TestUtils::create_new_policy_idx<3>(exec); + auto exec5 = TestUtils::create_new_policy_idx<4>(exec); + auto exec6 = TestUtils::create_new_policy_idx<5>(exec); + auto exec7 = TestUtils::create_new_policy_idx<6>(exec); + auto exec8 = TestUtils::create_new_policy_idx<7>(exec); + auto exec9 = TestUtils::create_new_policy_idx<8>(exec); + auto exec10 = TestUtils::create_new_policy_idx<9>(exec); + auto exec11 = TestUtils::create_new_policy_idx<10>(exec); + auto exec12 = TestUtils::create_new_policy_idx<11>(exec); + auto exec13 = TestUtils::create_new_policy_idx<12>(exec); + + oneapi::dpl::counting_iterator counting(size_t{0}); + + const auto n = last - first; + + //Run the tests + auto get_expect = [n](auto exp) { + if constexpr (__reverses % 2 == 0) + { + return exp; + } + else + { + return std::make_reverse_iterator(exp + n); + } + }; + +# if _ONEDPL_DEBUG_SYCL + std::cout << input_descr << ":"; +# endif // _ONEDPL_DEBUG_SYCL + + if constexpr (__read) + { + oneapi::dpl::fill(exec1, orig_out_first, orig_out_first + n, trash); + if constexpr (__reset_read) + { + //Reset data if required + oneapi::dpl::copy(exec2, expected_first, expected_first + n, orig_first); + } + + //Run test + oneapi::dpl::copy(exec3, first, last, copy_to_first); + + //get expected sequence with proper number of reverses + auto expect = get_expect(expected_first); + std::string msg = std::string("wrong read effect from ") + input_descr; + //verify result using original unwrapped output + EXPECT_EQ_N(expect, orig_out_first, n, msg.c_str()); +# if _ONEDPL_DEBUG_SYCL + std::cout << " read pass,"; +# endif // _ONEDPL_DEBUG_SYCL + } + if constexpr (__write) + { + //Reset data + if constexpr (__check_write) + { + //only reset output data if we intend to check it afterward + oneapi::dpl::fill(exec4, orig_first, orig_first + n, trash); + } + + oneapi::dpl::copy(exec5, copy_from_first, copy_from_first + n, first); + //check write if required (ignore discard iterator) + if constexpr (__check_write) + { + //copy back data from original unwrapped sequence + std::vector copy_back(n); + oneapi::dpl::copy(exec6, orig_first, orig_first + n, copy_back.begin()); + + //get expected sequence with proper number of reverses + auto expect = get_expect(expected_first); + std::string msg = std::string("wrong write effect from ") + input_descr; + //verify copied back data + EXPECT_EQ_N(expect, copy_back.begin(), n, msg.c_str()); +# if _ONEDPL_DEBUG_SYCL + std::cout << " write pass"; +# endif // _ONEDPL_DEBUG_SYCL + } + else + { +# if _ONEDPL_DEBUG_SYCL + std::cout << " write pass (no check)"; +# endif // _ONEDPL_DEBUG_SYCL + } + } + if constexpr (!__read && !__write) + { +# if _ONEDPL_DEBUG_SYCL + std::cout << " has no valid tests"; +# endif // _ONEDPL_DEBUG_SYCL + } +# if _ONEDPL_DEBUG_SYCL + std::cout << std::endl; +# endif // _ONEDPL_DEBUG_SYCL + + // Now recurse with a layer of wrappers if requested + if constexpr (__recurse > 0) + { +# if _ONEDPL_DEBUG_SYCL + std::cout << std::endl << "Recursing on " << input_descr << ":" << std::endl; +# endif // _ONEDPL_DEBUG_SYCL + oneapi::dpl::discard_iterator discard{}; + // iterate through all wrappers and recurse - 1 + auto noop = [](auto i) { return i; }; + + if constexpr (__is_reversible) + { // std::reverse_iterator(it) + auto reversed_first = ::std::make_reverse_iterator(last); + auto reversed_last = ::std::make_reverse_iterator(first); + std::string new_input_descr = std::string("std::reverse(") + input_descr + std::string(")"); + //TODO: Look at device copyability of std::reverse_iterator and re-enable recurse + wrap_recurse<0, __reverses + 1, __read, __reset_read, __write, __check_write, __usable_as_perm_map, + __usable_as_perm_src, __is_reversible>(exec7, reversed_first, reversed_last, copy_from_first, + copy_to_first, orig_first, orig_out_first, + expected_first, trash, new_input_descr); + } + + { //transform_iterator(it,noop{}) + auto trans = oneapi::dpl::make_transform_iterator(first, noop); + std::string new_input_descr = std::string("transform_iterator(") + input_descr + std::string(", noop)"); + wrap_recurse<__recurse - 1, __reverses, __read, __reset_read, /*__write=*/false, __check_write, + __usable_as_perm_map, __usable_as_perm_src, __is_reversible>( + exec8, trans, trans + n, discard, copy_to_first, orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + + if constexpr (__usable_as_perm_src) + { //permutation_iteartor(it,noop{}) + std::string new_input_descr = std::string("permutation_iterator(") + input_descr + std::string(", noop)"); + auto perm = oneapi::dpl::make_permutation_iterator(first, noop); + wrap_recurse<__recurse - 1, __reverses, __read, __reset_read, __write, __check_write, __usable_as_perm_map, + __usable_as_perm_src, __is_reversible>(exec9, perm, perm + n, copy_from_first, copy_to_first, + orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + + if constexpr (__usable_as_perm_src) + { //permutation_iterator(it,counting_iter) + std::string new_input_descr = + std::string("permutation_iterator(") + input_descr + std::string(", counting_iterator)"); + auto perm = oneapi::dpl::make_permutation_iterator(first, counting); + wrap_recurse<__recurse - 1, __reverses, __read, __reset_read, __write, __check_write, __usable_as_perm_map, + __usable_as_perm_src, __is_reversible>(exec11, perm, perm + n, copy_from_first, copy_to_first, + orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + + if constexpr (__usable_as_perm_map) + { //permutation_iterator(counting_iterator,it) + std::string new_input_descr = + std::string("permutation_iterator(counting_iterator,") + input_descr + std::string(")"); + auto perm = oneapi::dpl::make_permutation_iterator(counting, first); + wrap_recurse<__recurse - 1, __reverses, __read, __reset_read, /*__write=*/false, __check_write, + __usable_as_perm_map, __usable_as_perm_src, __is_reversible>( + exec10, perm, perm + n, discard, copy_to_first, orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + + { //zip_iterator(counting_iterator,it) + std::string new_input_descr = + std::string("zip_iterator(counting_iterator,") + input_descr + std::string(")"); + auto zip = oneapi::dpl::make_zip_iterator(counting, first); + auto zip_out = oneapi::dpl::make_zip_iterator(discard, copy_to_first); + wrap_recurse<__recurse - 1, __reverses, __read, __reset_read, /*__write=*/false, __check_write, + /*__usable_as_perm_map=*/false, __usable_as_perm_src, __is_reversible>( + exec12, zip, zip + n, discard, zip_out, orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + + { //zip_iterator(it, discard_iterator) + std::string new_input_descr = + std::string("zip_iterator(") + input_descr + std::string(", discard_iterator)"); + auto zip = oneapi::dpl::make_zip_iterator(first, discard); + auto zip_in = oneapi::dpl::make_zip_iterator(copy_from_first, counting); + wrap_recurse<__recurse - 1, __reverses, /*__read=*/false, false, __write, __check_write, + /*__usable_as_perm_map=*/false, __usable_as_perm_src, __is_reversible>( + exec13, zip, zip + n, zip_in, discard, orig_first, orig_out_first, expected_first, trash, + new_input_descr); + } + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +#endif //_INPUT_DATA_SWEEP_H diff --git a/test/parallel_api/iterator/input_data_sweep_counting_iter.pass.cpp b/test/parallel_api/iterator/input_data_sweep_counting_iter.pass.cpp new file mode 100644 index 00000000000..da9bb5d7977 --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_counting_iter.pass.cpp @@ -0,0 +1,94 @@ +// -*- C++ -*- +//===-- input_data_sweep_counting_iter.pass.cpp ---------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +#if TEST_DPCPP_BACKEND_PRESENT + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +template +void +test(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if constexpr (std::is_integral_v) + { + if (TestUtils::has_types_support(policy.queue().get_device())) + { + + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + oneapi::dpl::counting_iterator my_counting(0); + //counting_iterator + wrap_recurse<__recurse, 0, /*__read =*/true, /*__reset_read=*/false, /*__write=*/false, + /*__check_write=*/false, /*__usable_as_perm_map=*/true, /*__usable_as_perm_src=*/true, + /*__is_reversible=*/true>(policy, my_counting, my_counting + n, counting, copy_out.get_data(), + my_counting, copy_out.get_data(), counting, trash, + std::string("counting_iterator<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + auto policy5 = TestUtils::create_new_policy_idx<4>(policy); + + // baseline with no wrapping + test(policy1, -666.0f, n, "float"); + test(policy2, -666.0, n, "double"); + test(policy3, 999, n, "uint64_t"); + + // big recursion step: 1 and 2 layers of wrapping + test(policy4, -666, n, "int32_t"); + + // special case: discard iterator + oneapi::dpl::counting_iterator counting(0); + oneapi::dpl::discard_iterator discard{}; + wrap_recurse<1, 0, /*__read =*/false, /*__reset_read=*/false, /*__write=*/true, + /*__check_write=*/false, /*__usable_as_perm_map=*/false, /*__usable_as_perm_src=*/true, + /*__is_reversible=*/true>(policy5, discard, discard + n, counting, discard, discard, discard, discard, + -666, "discard_iterator"); + +#endif // TEST_DPCPP_BACKEND_PRESENT + + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/parallel_api/iterator/input_data_sweep_host_iter.pass.cpp b/test/parallel_api/iterator/input_data_sweep_host_iter.pass.cpp new file mode 100644 index 00000000000..65809b90ffa --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_host_iter.pass.cpp @@ -0,0 +1,82 @@ +// -*- C++ -*- +//===-- input_data_sweep_host_iter.pass.cpp -------------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +#if TEST_DPCPP_BACKEND_PRESENT + +template +void +test(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + auto copy_from = oneapi::dpl::counting_iterator(0); + // host iterator + std::vector host_iter(n); + wrap_recurse<__recurse, 0, /*__read =*/true, /*__reset_read=*/true, /*__write=*/true, + /*__check_write=*/true, /*__usable_as_perm_map=*/true, /*__usable_as_perm_src=*/false, + /*__is_reversible=*/true>(policy, host_iter.begin(), host_iter.end(), copy_from, + copy_out.get_data(), host_iter.begin(), copy_out.get_data(), copy_from, + trash, std::string("host_iterator<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + + // baseline with no wrapping + test(policy1, -666.0f, n, "float"); + test(policy2, -666.0, n, "double"); + test(policy3, 999, n, "uint64_t"); + + // big recursion step: 1 and 2 layers of wrapping + test(policy4, -666, n, "int32_t"); + +#endif // TEST_DPCPP_BACKEND_PRESENT + + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/parallel_api/iterator/input_data_sweep_sycl_iter.pass.cpp b/test/parallel_api/iterator/input_data_sweep_sycl_iter.pass.cpp new file mode 100644 index 00000000000..d612ddb461c --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_sycl_iter.pass.cpp @@ -0,0 +1,83 @@ +// -*- C++ -*- +//===-- input_data_sweep_sycl_iter.pass.cpp -------------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +#if TEST_DPCPP_BACKEND_PRESENT + +template +void +test(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + // sycl iterator + sycl::buffer buf(n); + //test all modes / wrappers + wrap_recurse<__recurse, 0, /*__read =*/true, /*__reset_read=*/true, /*__write=*/true, + /*__check_write=*/true, /*__usable_as_perm_map=*/true, /*__usable_as_perm_src=*/true, + /*__is_reversible=*/false>(policy, oneapi::dpl::begin(buf), oneapi::dpl::end(buf), counting, + copy_out.get_data(), oneapi::dpl::begin(buf), copy_out.get_data(), + counting, trash, + std::string("sycl_iterator<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + + // baseline with no wrapping + test(policy1, -666.0f, n, "float"); + test(policy2, -666.0, n, "double"); + test(policy3, 999, n, "uint64_t"); + + // big recursion step: 1 and 2 layers of wrapping + test(policy4, -666, n, "int32_t"); + +#endif // TEST_DPCPP_BACKEND_PRESENT + + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/parallel_api/iterator/input_data_sweep_usm_allocator.pass.cpp b/test/parallel_api/iterator/input_data_sweep_usm_allocator.pass.cpp new file mode 100644 index 00000000000..90b635a4e86 --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_usm_allocator.pass.cpp @@ -0,0 +1,119 @@ +// -*- C++ -*- +//===-- input_data_sweep_usm_allocator.pass.cpp ---------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +#if TEST_DPCPP_BACKEND_PRESENT + +template +void +test_usm_shared_alloc(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + //std::vector using usm shared allocator + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + // usm_shared allocator std::vector + sycl::usm_allocator q_alloc{policy.queue()}; + std::vector shared_data_vec(n, q_alloc); + //test all modes / wrappers + + //Only test as source iterator for permutation iterator if we can expect it to work + // (if the vector implementation distiguishes its iterator for this type) + wrap_recurse< + __recurse, 0, /*__read =*/true, /*__reset_read=*/true, /*__write=*/true, + /*__check_write=*/true, /*__usable_as_perm_map=*/true, + /*__usable_as_perm_src=*/ + TestUtils::__vector_impl_distinguishes_usm_allocator_from_default::value, + /*__is_reversible=*/true>(policy, shared_data_vec.begin(), shared_data_vec.end(), counting, + copy_out.get_data(), shared_data_vec.begin(), copy_out.get_data(), counting, + trash, std::string("usm_shared_alloc_vector<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} + +template +void +test_usm_host_alloc(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + //std::vector using usm host allocator + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + // usm_host allocator std::vector + sycl::usm_allocator q_alloc{policy.queue()}; + std::vector host_data_vec(n, q_alloc); + //test all modes / wrappers + + //Only test as source iterator for permutation iterator if we can expect it to work + // (if the vector implementation distiguishes its iterator for this type) + wrap_recurse< + __recurse, 0, /*__read =*/true, /*__reset_read=*/true, /*__write=*/true, + /*__check_write=*/true, /*__usable_as_perm_map=*/true, /*__usable_as_perm_src=*/ + TestUtils::__vector_impl_distinguishes_usm_allocator_from_default::value, + /*__is_reversible=*/true>(policy, host_data_vec.begin(), host_data_vec.end(), counting, copy_out.get_data(), + host_data_vec.begin(), copy_out.get_data(), counting, trash, + std::string("usm_host_alloc_vector<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + auto policy5 = TestUtils::create_new_policy_idx<4>(policy); + + // baseline with no wrapping + test_usm_shared_alloc(policy1, -666.0f, n, "float"); + test_usm_shared_alloc(policy2, -666.0, n, "double"); + test_usm_shared_alloc(policy3, 999, n, "uint64_t"); + // big recursion step: 1 and 2 layers of wrapping + test_usm_shared_alloc(policy4, -666, n, "int32_t"); + //only use host alloc for int, it follows the same path as shared alloc + test_usm_host_alloc(policy5, 666, n, "int"); + +#endif // TEST_DPCPP_BACKEND_PRESENT + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/parallel_api/iterator/input_data_sweep_usm_device.pass.cpp b/test/parallel_api/iterator/input_data_sweep_usm_device.pass.cpp new file mode 100644 index 00000000000..4c08fb8f6ad --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_usm_device.pass.cpp @@ -0,0 +1,81 @@ +// -*- C++ -*- +//===-- input_data_sweep_usm_device.pass.cpp ------------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +#if TEST_DPCPP_BACKEND_PRESENT + +template +void +test(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + // usm_device + TestUtils::usm_data_transfer device_data(policy.queue(), n); + auto usm_device = device_data.get_data(); + //test all modes / wrappers + wrap_recurse<__recurse, 0>(policy, usm_device, usm_device + n, counting, copy_out.get_data(), usm_device, + copy_out.get_data(), counting, trash, + std::string("usm_device<") + type_text + std::string(">")); + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + + // baseline with no wrapping + test(policy1, -666.0f, n, "float"); + test(policy2, -666.0, n, "double"); + test(policy3, 999, n, "uint64_t"); + + // big recursion step: 1 and 2 layers of wrapping + test(policy4, -666, n, "int32_t"); + +#endif // TEST_DPCPP_BACKEND_PRESENT + + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/parallel_api/iterator/input_data_sweep_usm_shared.pass.cpp b/test/parallel_api/iterator/input_data_sweep_usm_shared.pass.cpp new file mode 100644 index 00000000000..b1031804b03 --- /dev/null +++ b/test/parallel_api/iterator/input_data_sweep_usm_shared.pass.cpp @@ -0,0 +1,96 @@ +// -*- C++ -*- +//===-- input_data_sweep_usm_shared.pass.cpp ------------------------------===// +// +// Copyright (C) Intel Corporation +// +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// This file incorporates work covered by the following copyright and permission +// notice: +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// +//===----------------------------------------------------------------------===// + +#include "support/utils.h" +#include _PSTL_TEST_HEADER(execution) +#include _PSTL_TEST_HEADER(algorithm) +#include _PSTL_TEST_HEADER(iterator) + +#include "input_data_sweep.h" + +#include "support/utils_invoke.h" + +//This test is written without indirection from invoke_on_all_hetero_policies to make clear exactly which types +// are being tested, and to limit the number of types to be within reason. + +#if TEST_DPCPP_BACKEND_PRESENT + +template +void +test(Policy&& policy, T trash, size_t n, const std::string& type_text) +{ + if (TestUtils::has_types_support(policy.queue().get_device())) + { + + { //usm shared ptr + TestUtils::usm_data_transfer copy_out(policy.queue(), n); + oneapi::dpl::counting_iterator counting(0); + // usm_shared + TestUtils::usm_data_transfer shared_data(policy.queue(), n); + auto usm_shared = shared_data.get_data(); + //test all modes / wrappers + wrap_recurse<__recurse, 0>(policy, usm_shared, usm_shared + n, counting, copy_out.get_data(), usm_shared, + copy_out.get_data(), counting, trash, + std::string("usm_shared<") + type_text + std::string(">")); + } + } + else + { + TestUtils::unsupported_types_notifier(policy.queue().get_device()); + } +} + +#endif //TEST_DPCPP_BACKEND_PRESENT + +int +main() +{ +#if TEST_DPCPP_BACKEND_PRESENT + + constexpr size_t n = 10; + + auto q = TestUtils::get_test_queue(); + + auto policy = TestUtils::make_new_policy(q); + + auto policy1 = TestUtils::create_new_policy_idx<0>(policy); + auto policy2 = TestUtils::create_new_policy_idx<1>(policy); + auto policy3 = TestUtils::create_new_policy_idx<2>(policy); + auto policy4 = TestUtils::create_new_policy_idx<3>(policy); + auto policy5 = TestUtils::create_new_policy_idx<4>(policy); + + // baseline with no wrapping + test(policy1, -666.0f, n, "float"); + test(policy2, -666.0, n, "double"); + test(policy3, 999, n, "uint64_t"); + + // big recursion step: 1 and 2 layers of wrapping + test(policy4, -666, n, "int32_t"); + + // special case: recurse once on perm(perm(usm_shared,count), count) + oneapi::dpl::counting_iterator counting(0); + TestUtils::usm_data_transfer copy_out(policy5.queue(), n); + TestUtils::usm_data_transfer input(policy5.queue(), n); + auto perm1 = oneapi::dpl::make_permutation_iterator(input.get_data(), counting); + auto perm2 = oneapi::dpl::make_permutation_iterator(perm1, counting); + wrap_recurse<1, 0, /*__read =*/false, /*__reset_read=*/false, /*__write=*/true, + /*__check_write=*/false, /*__usable_as_perm_map=*/true, /*__usable_as_perm_src=*/true, + /*__is_reversible=*/true>( + policy5, perm2, perm2 + n, counting, copy_out.get_data(), perm2, copy_out.get_data(), counting, -666, + "permutation_iter(permutation_iterator(usm_shared,counting_iterator),counting_iterator)"); +#endif // TEST_DPCPP_BACKEND_PRESENT + + return TestUtils::done(TEST_DPCPP_BACKEND_PRESENT); +} diff --git a/test/support/utils_invoke.h b/test/support/utils_invoke.h index 2c66968f255..90ea9cf8966 100644 --- a/test/support/utils_invoke.h +++ b/test/support/utils_invoke.h @@ -75,6 +75,40 @@ make_fpga_policy(Arg&& arg) } #endif // _ONEDPL_FPGA_DEVICE +//function is needed to wrap kernel name into another class +template = 0> +auto +make_new_policy(_Policy&& __policy) + -> decltype(TestUtils::make_device_policy<_NewKernelName>(::std::forward<_Policy>(__policy))) +{ + return TestUtils::make_device_policy<_NewKernelName>(::std::forward<_Policy>(__policy)); +} + +#if ONEDPL_FPGA_DEVICE +template = 0> +auto +make_new_policy(_Policy&& __policy) + -> decltype(TestUtils::make_fpga_policy<::std::decay_t<_Policy>::unroll_factor, _NewKernelName>( + ::std::forward<_Policy>(__policy))) +{ + return TestUtils::make_fpga_policy<::std::decay_t<_Policy>::unroll_factor, _NewKernelName>( + ::std::forward<_Policy>(__policy)); +} +#endif + +template +auto +make_new_policy(sycl::queue _queue) +{ +#if ONEDPL_FPGA_DEVICE + return TestUtils::make_fpga_policy(_queue); +#else + return TestUtils::make_device_policy(_queue); +#endif +} + #endif // TEST_DPCPP_BACKEND_PRESENT //////////////////////////////////////////////////////////////////////////////// @@ -182,12 +216,7 @@ struct invoke_on_all_hetero_policies // performs some checks that fail. As a workaround, define for functors which have this issue // __functor_type(see kernel_type definition) type field which doesn't have any pointers in it's name. using kernel_name = unique_kernel_name; - auto my_policy = -#if ONEDPL_FPGA_DEVICE - TestUtils::make_fpga_policy(queue); -#else - TestUtils::make_device_policy(queue); -#endif + auto my_policy = make_new_policy(queue); iterator_invoker<::std::random_access_iterator_tag, /*IsReverse*/ ::std::false_type>()( my_policy, op, ::std::forward(rest)...); } diff --git a/test/support/utils_sycl.h b/test/support/utils_sycl.h index 9acb71f2650..ad17ae71be3 100644 --- a/test/support/utils_sycl.h +++ b/test/support/utils_sycl.h @@ -75,29 +75,6 @@ auto async_handler = [](sycl::exception_list ex_list) { } }; -//function is needed to wrap kernel name into another class -template = 0> -auto -make_new_policy(_Policy&& __policy) - -> decltype(TestUtils::make_device_policy<_NewKernelName>(::std::forward<_Policy>(__policy))) -{ - return TestUtils::make_device_policy<_NewKernelName>(::std::forward<_Policy>(__policy)); -} - -#if ONEDPL_FPGA_DEVICE -template = 0> -auto -make_new_policy(_Policy&& __policy) - -> decltype(TestUtils::make_fpga_policy<::std::decay_t<_Policy>::unroll_factor, _NewKernelName>( - ::std::forward<_Policy>(__policy))) -{ - return TestUtils::make_fpga_policy<::std::decay_t<_Policy>::unroll_factor, _NewKernelName>( - ::std::forward<_Policy>(__policy)); -} -#endif - #if ONEDPL_FPGA_DEVICE inline auto default_selector = # if ONEDPL_FPGA_EMULATOR