diff options
author | MikeDvorskiy <40659946+MikeDvorskiy@users.noreply.github.com> | 2021-03-04 12:49:20 +0300 |
---|---|---|
committer | GitHub <noreply@github.com> | 2021-03-04 12:49:20 +0300 |
commit | 0b847654bd4f13aad20ec839822e8a82b506e0bf (patch) | |
tree | 57b628aa4d20c88f034bd63639f4d92dc6cea9d0 | |
parent | [pstl][host] + a hot fix for "swap" algorithm; (passing is_parallel tag) (diff) | |
download | llvm-project-0b847654bd4f13aad20ec839822e8a82b506e0bf.tar.gz llvm-project-0b847654bd4f13aad20ec839822e8a82b506e0bf.tar.bz2 llvm-project-0b847654bd4f13aad20ec839822e8a82b506e0bf.zip |
Ranges API support: passing "sycl buffer" as "pseudo-range" (#121)
* [dpc++][ranges] changes in CPO "views::all" - now it can accept any range
* [dpc++][ranges] + passing "sycl buffer" as "pseudo-range"
* [dpc++][ranges] the tests for passing "sycl buffer" as "pseudo-range"
* [dpc++][ranges] + fixes for "passing sycl buffer as pseudo-range"
* [dpc++] removed unused "stuff" from dpcpp(SYCL) backend
* Extend CI tests (#110)
* Extend CI tests
* Added 8 to -j make parameter
* Updated envs path for tests
* Deleted Tests_backend_tbb_cxx_11 test from Windows check
* Deleted Tests_backend_tbb_cxx_17 test from Windows check
* Updated tests description
* Remove unnecessary barriers and add constexpr attribute in scan brick (#119)
* Add constexpr attribute to__iters_per_witem variable
Signed-off-by: Pavlov, Evgeniy <evgeniy.pavlov@intel.com>
* Remove unnesessary barriers
Signed-off-by: Pavlov, Evgeniy <evgeniy.pavlov@intel.com>
* Remove auto to decltype
Signed-off-by: Pavlov, Evgeniy <evgeniy.pavlov@intel.com>
* Remove extra class from parallel_backend_sycl.h (#120)
Signed-off-by: Sobolev, Dmitriy <dmitriy.sobolev@intel.com>
* [pstl] + a hot fix for an identity value for __pattern_minmax_element
* Add FPGA_EMU tests (#118)
* Refactor generated names of kernels (#113)
* reworked kernel names
* more testing is added
* make shift_left_right run with unnamed lambdas
* some fixes for CI
* applied some suggestions from review
* addressed feedback from review
* cmake changes + kernel name for shift_left, shift_right algorithms
* make some tests compile
* remove unused variables
* Update theme and index files (#124)
* CMake: add openmp-simd compiler option (#116)
* Fix for reduction stage of scan pattern (#125)
* fixed debug issue
* some cosmetic changes
* Update comment
Co-authored-by: Dmitriy Sobolev <Dmitriy.Sobolev@intel.com>
* remove extra braces
Co-authored-by: Dmitriy Sobolev <Dmitriy.Sobolev@intel.com>
* Add _PSTL_PRAGMA_SIMD_ORDERED_MONOTONIC_2ARGS definition in case of the use of standard library headers (#127)
Signed-off-by: Pavlov, Evgeniy <evgeniy.pavlov@intel.com>
* Remove explicit default copy constructor in copy_constructible_value_holder (#123)
* CTest: skipt test if return code is 77
* Supposed solution
* Erased explicit copy c'tor in copy_assignable_value_holder, added is_default_constructible check in iterators test
* Deleted extra changes
* Deleted extra changes
* Splited static_assert in two lines: bool expression and message
* Splited assert in two lines
* Use shorter static_assert version with using
* Inserted spaces where it needs
* Reverted test calls and clarified commentaries
* Avoid divergence of work items in the same SIMD before calling collectives (#129)
* Fix for kernel names check (#133)
* Refactor generated names of kernels (#113)
* reworked kernel names
* more testing is added
* make shift_left_right run with unnamed lambdas
* some fixes for CI
* applied some suggestions from review
* addressed feedback from review
* cmake changes + kernel name for shift_left, shift_right algorithms
* make some tests compile
* remove unused variables
* Version update (#132)
* Fix segmentation fault in the global scan functor (#139)
* Fix a check in the global scan functor
Signed-off-by: Pavlov, Evgeniy <evgeniy.pavlov@intel.com>
* Update include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h
Co-authored-by: Dmitriy Sobolev <Dmitriy.Sobolev@intel.com>
Co-authored-by: Andrey Fedorov <andrey.fedorov@intel.com>
Co-authored-by: Dmitriy Sobolev <Dmitriy.Sobolev@intel.com>
* Turn off strict aliasing optimization if used (#137)
* Turn off strict aliasing optimization when use icpc on Linux
* Fix usage of CMAKE_CXX_COMPILER and CMAKE_BUILD_TYPE variables
* Make changes affect only tests
* Get rid of nested if condition
* Async api extensions (#78)
* adding async API as an experimental feature
* initial support for DPCPP backend only
* implementation for subset of algorithm/numeric (copy,fill,for_each,sort,reduce,transform,transform_reduce) with suffix async; accepting an arbitrary number of sycl::event's as last argument to express input dependencies
* returning a future-like object of undefined type that is convertible into a sycl::event.
* lifetime of temporary storage is bound to lifetime of returned object
* [dpc++][ranges] the tests for passing "sycl buffer" as "pseudo-range"
* Reduce number of elements to scan due to taking too much time when testing (#142)
* Reduce number of elements to scan because of taking too much time
* Do not change number of elements when using release mode
Signed-off-by: Sobolev, Dmitriy <dmitriy.sobolev@intel.com>
* [dpc++][ranges] the tests: making unique kernel name
* [dpc++][ranges] find_sycl_range_pass.cpp test: making unique kernel name
* [dpc++][ranges] + minor improvements in meta-programming (passing "sycl buffer" as "pseudo-range")
* + clang format
* + decay
* temporary checking a negative test
* Revert "temporary checking a negative test"
This reverts commit 8ced2ef374eaf23d82572add3c8400e047d279c6.
Co-authored-by: Alexey Oralov <alexey.oralov@intel.com>
Co-authored-by: Evgeniy Pavlov <evgeniy.pavlov@intel.com>
Co-authored-by: Dmitriy Sobolev <Dmitriy.Sobolev@intel.com>
Co-authored-by: Andrey Fedorov <andrey.fedorov@intel.com>
Co-authored-by: Valentina Kats <valentina.kats@intel.com>
Co-authored-by: Alexey Veprev <alexey.veprev@intel.com>
Co-authored-by: Denis Paranichev <denis.paranichev@intel.com>
Co-authored-by: Pablo Reble <pablo@reble.org>
27 files changed, 141 insertions, 93 deletions
diff --git a/include/oneapi/dpl/pstl/execution_defs.h b/include/oneapi/dpl/pstl/execution_defs.h index d6c2fed8a5fd..a8225ea4c187 100644 --- a/include/oneapi/dpl/pstl/execution_defs.h +++ b/include/oneapi/dpl/pstl/execution_defs.h @@ -201,13 +201,30 @@ template <typename _ExecPolicy, typename _T> using __ref_or_copy = typename oneapi::dpl::__internal::__ref_or_copy_impl<typename ::std::decay<_ExecPolicy>::type, _T>::type; +// utilities for Range API template <typename _R> -using __difference_t = decltype(::std::declval<_R&>().size()); +auto +__check_size(int) -> decltype(::std::declval<_R&>().size()); + +template <typename _R> +auto +__check_size(...) -> decltype(::std::declval<_R&>().get_count()); + +template <typename _R> +using __difference_t = decltype(__check_size<_R>(0)); + +template <typename _R> +auto +__check_subscript(int) -> typename ::std::decay<decltype(::std::declval<_R&>()[0])>::type; + +template <typename _R> +auto +__check_subscript(...) -> typename ::std::decay<_R>::type::value_type; template <typename _R> struct __range_traits { - using __value_t = typename ::std::decay<decltype(::std::declval<_R&>()[0])>::type; + using __value_t = decltype(__check_subscript<_R>(0)); }; template <typename _R> diff --git a/include/oneapi/dpl/pstl/glue_algorithm_ranges_impl.h b/include/oneapi/dpl/pstl/glue_algorithm_ranges_impl.h index b3af016103c7..46ca4a9454f8 100644 --- a/include/oneapi/dpl/pstl/glue_algorithm_ranges_impl.h +++ b/include/oneapi/dpl/pstl/glue_algorithm_ranges_impl.h @@ -40,7 +40,7 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> for_each(_ExecutionPolicy&& __exec, _Range&& __rng, _Function __f) { oneapi::dpl::__internal::__ranges::__pattern_walk1(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), __f); + views::all(::std::forward<_Range>(__rng)), __f); } // [alg.find] @@ -50,7 +50,7 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi:: find_if(_ExecutionPolicy&& __exec, _Range&& __rng, _Predicate __pred) { return oneapi::dpl::__internal::__ranges::__pattern_find_if(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), __pred); + views::all_read(::std::forward<_Range>(__rng)), __pred); } template <typename _ExecutionPolicy, typename _Range, typename _Predicate> @@ -80,9 +80,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range1>> find_end(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _BinaryPredicate __pred) { - return oneapi::dpl::__internal::__ranges::__pattern_find_end(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range1>(__rng1), - ::std::forward<_Range2>(__rng2), __pred); + return oneapi::dpl::__internal::__ranges::__pattern_find_end( + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), __pred); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2> @@ -101,9 +101,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range1>> find_first_of(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _BinaryPredicate __pred) { - return oneapi::dpl::__internal::__ranges::__pattern_find_first_of(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range1>(__rng1), - ::std::forward<_Range2>(__rng2), __pred); + return oneapi::dpl::__internal::__ranges::__pattern_find_first_of( + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), __pred); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2> @@ -122,9 +122,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range1>> search(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _BinaryPredicate __pred) { - return oneapi::dpl::__internal::__ranges::__pattern_search(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range1>(__rng1), - ::std::forward<_Range2>(__rng2), __pred); + return oneapi::dpl::__internal::__ranges::__pattern_search( + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), __pred); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2> @@ -140,8 +140,9 @@ template <typename _ExecutionPolicy, typename _Range, typename _Size, typename _ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range>> search_n(_ExecutionPolicy&& __exec, _Range&& __rng, _Size __count, const _Tp& __value, _BinaryPredicate __pred) { - return oneapi::dpl::__internal::__ranges::__pattern_search_n( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range>(__rng), __count, __value, __pred); + return oneapi::dpl::__internal::__ranges::__pattern_search_n(::std::forward<_ExecutionPolicy>(__exec), + views::all_read(::std::forward<_Range>(__rng)), + __count, __value, __pred); } template <typename _ExecutionPolicy, typename _Range, typename _Size, typename _Tp> @@ -159,8 +160,8 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> copy(_ExecutionPolicy&& __exec, _Range1&& __rng, _Range2&& __result) { oneapi::dpl::__internal::__ranges::__pattern_walk2( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng), ::std::forward<_Range2>(__result), - oneapi::dpl::__internal::__brick_copy<_ExecutionPolicy>{}); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng)), + views::all_write(::std::forward<_Range2>(__result)), oneapi::dpl::__internal::__brick_copy<_ExecutionPolicy>{}); } // [alg.transform] @@ -170,8 +171,8 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> transform(_ExecutionPolicy&& __exec, _Range1&& __rng, _Range2&& __result, _UnaryOperation __op) { oneapi::dpl::__internal::__ranges::__pattern_walk2( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng), ::std::forward<_Range2>(__result), - [__op](auto x, auto& z) { z = __op(x); }); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng)), + views::all_write(::std::forward<_Range2>(__result)), [__op](auto x, auto& z) { z = __op(x); }); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3, typename _BinaryOperation> @@ -179,8 +180,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> transform(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __result, _BinaryOperation __op) { oneapi::dpl::__internal::__ranges::__pattern_walk3( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - ::std::forward<_Range3>(__result), [__op](auto x, auto y, auto& z) { z = __op(x, y); }); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), views::all_write(::std::forward<_Range3>(__result)), + [__op](auto x, auto y, auto& z) { z = __op(x, y); }); } // [alg.sort] @@ -190,16 +192,15 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> sort(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp) { oneapi::dpl::__internal::__ranges::__pattern_sort(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), __comp); + views::all(::std::forward<_Range>(__rng)), __comp); } template <typename _ExecutionPolicy, typename _Range> oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, void> sort(_ExecutionPolicy&& __exec, _Range&& __rng) { - oneapi::dpl::__internal::__ranges::__pattern_sort(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), - oneapi::dpl::__internal::__pstl_less()); + sort(::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range>(__rng), + oneapi::dpl::__internal::__pstl_less()); } // [stable.sort] @@ -225,11 +226,12 @@ template <typename _ExecutionPolicy, typename _Range, typename _Compare> oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range>> is_sorted_until(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp) { + auto __view = views::all_read(::std::forward<_Range>(__rng)); const auto __res = oneapi::dpl::__internal::__ranges::__pattern_adjacent_find( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range>(__rng), - oneapi::dpl::__internal::__reorder_pred<_Compare>(__comp), oneapi::dpl::__internal::__first_semantic()); + ::std::forward<_ExecutionPolicy>(__exec), __view, oneapi::dpl::__internal::__reorder_pred<_Compare>(__comp), + oneapi::dpl::__internal::__first_semantic()); - return __res == __rng.size() ? __res : __res + 1; + return __res == __view.size() ? __res : __res + 1; } template <typename _ExecutionPolicy, typename _Range> @@ -244,10 +246,11 @@ template <typename _ExecutionPolicy, typename _Range, typename _Compare> oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, bool> is_sorted(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp) { + auto __view = views::all_read(::std::forward<_Range>(__rng)); return oneapi::dpl::__internal::__ranges::__pattern_adjacent_find( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range>(__rng), + ::std::forward<_ExecutionPolicy>(__exec), __view, oneapi::dpl::__internal::__reorder_pred<_Compare>(__comp), - oneapi::dpl::__internal::__or_semantic()) == __rng.size(); + oneapi::dpl::__internal::__or_semantic()) == __view.size(); } template <typename _ExecutionPolicy, typename _Range> @@ -266,8 +269,8 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, merge(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __rng3, _Compare __comp) { return oneapi::dpl::__internal::__ranges::__pattern_merge( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - ::std::forward<_Range3>(__rng3), __comp); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), views::all_write(::std::forward<_Range3>(__rng3)), __comp); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Range3> @@ -275,9 +278,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range3>> merge(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _Range3&& __rng3) { - return oneapi::dpl::__internal::__ranges::__pattern_merge( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - ::std::forward<_Range3>(__rng3), oneapi::dpl::__internal::__pstl_less()); + return merge(::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), + ::std::forward<_Range2>(__rng2), ::std::forward<_Range3>(__rng3), + oneapi::dpl::__internal::__pstl_less()); } // [alg.min.max] @@ -286,8 +289,8 @@ template <typename _ExecutionPolicy, typename _Range, typename _Compare> oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, oneapi::dpl::__internal::__difference_t<_Range>> min_element(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp) { - return oneapi::dpl::__internal::__ranges::__pattern_min_element(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), __comp); + return oneapi::dpl::__internal::__ranges::__pattern_min_element( + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range>(__rng)), __comp); } template <typename _ExecutionPolicy, typename _Range> @@ -321,8 +324,8 @@ oneapi::dpl::__internal::__enable_if_execution_policy< ::std::pair<oneapi::dpl::__internal::__difference_t<_Range>, oneapi::dpl::__internal::__difference_t<_Range>>> minmax_element(_ExecutionPolicy&& __exec, _Range&& __rng, _Compare __comp) { - return oneapi::dpl::__internal::__ranges::__pattern_minmax_element(::std::forward<_ExecutionPolicy>(__exec), - ::std::forward<_Range>(__rng), __comp); + return oneapi::dpl::__internal::__ranges::__pattern_minmax_element( + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range>(__rng)), __comp); } template <typename _ExecutionPolicy, typename _Range> diff --git a/include/oneapi/dpl/pstl/glue_numeric_ranges_impl.h b/include/oneapi/dpl/pstl/glue_numeric_ranges_impl.h index 217f63239db3..b0663539b263 100644 --- a/include/oneapi/dpl/pstl/glue_numeric_ranges_impl.h +++ b/include/oneapi/dpl/pstl/glue_numeric_ranges_impl.h @@ -65,8 +65,9 @@ transform_reduce(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, { using _ValueType = oneapi::dpl::__internal::__value_t<_Range1>; return oneapi::dpl::__internal::__ranges::__pattern_transform_reduce( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - __init, ::std::plus<_ValueType>(), ::std::multiplies<_ValueType>()); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), __init, ::std::plus<_ValueType>(), + ::std::multiplies<_ValueType>()); } template <typename _ExecutionPolicy, typename _Range1, typename _Range2, typename _Tp, typename _BinaryOperation1, @@ -76,8 +77,8 @@ transform_reduce(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& __rng2, _BinaryOperation1 __binary_op1, _BinaryOperation2 __binary_op2) { return oneapi::dpl::__internal::__ranges::__pattern_transform_reduce( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - __init, __binary_op1, __binary_op2); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_read(::std::forward<_Range2>(__rng2)), __init, __binary_op1, __binary_op2); } template <typename _ExecutionPolicy, typename _Range, typename _Tp, typename _BinaryOperation, typename _UnaryOperation> @@ -85,8 +86,9 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, _Tp> transform_reduce(_ExecutionPolicy&& __exec, _Range&& __rng, _Tp __init, _BinaryOperation __binary_op, _UnaryOperation __unary_op) { - return oneapi::dpl::__internal::__ranges::__pattern_transform_reduce( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range>(__rng), __init, __binary_op, __unary_op); + return oneapi::dpl::__internal::__ranges::__pattern_transform_reduce(::std::forward<_ExecutionPolicy>(__exec), + views::all_read(::std::forward<_Range>(__rng)), + __init, __binary_op, __unary_op); } // [exclusive.scan] @@ -153,8 +155,9 @@ transform_exclusive_scan(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& _BinaryOperation __binary_op, _UnaryOperation __unary_op) { return oneapi::dpl::__internal::__ranges::__pattern_transform_scan( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - __unary_op, __init, __binary_op, /*inclusive=*/::std::false_type()); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_write(::std::forward<_Range2>(__rng2)), __unary_op, __init, __binary_op, + /*inclusive=*/::std::false_type()); } // [transform.inclusive.scan] @@ -167,8 +170,8 @@ transform_inclusive_scan(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& _UnaryOperation __unary_op, _Tp __init) { return oneapi::dpl::__internal::__ranges::__pattern_transform_scan( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - __unary_op, __init, __binary_op, + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_write(::std::forward<_Range2>(__rng2)), __unary_op, __init, __binary_op, /*inclusive=*/::std::true_type()); } @@ -180,9 +183,8 @@ transform_inclusive_scan(_ExecutionPolicy&& __exec, _Range1&& __rng1, _Range2&& _UnaryOperation __unary_op) { return oneapi::dpl::__internal::__ranges::__pattern_transform_scan( - ::std::forward<_ExecutionPolicy>(__exec), ::std::forward<_Range1>(__rng1), ::std::forward<_Range2>(__rng2), - __unary_op, __binary_op, - /*inclusive=*/::std::true_type()); + ::std::forward<_ExecutionPolicy>(__exec), views::all_read(::std::forward<_Range1>(__rng1)), + views::all_write(::std::forward<_Range2>(__rng2)), __unary_op, __binary_op, /*inclusive=*/::std::true_type()); } } // namespace ranges diff --git a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h index 7096d0a6c2c4..fbd62a77a5f4 100644 --- a/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h +++ b/include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h @@ -81,9 +81,10 @@ class all_view accessor_t m_acc; }; +template <sycl::access::mode AccMode = sycl::access::mode::read_write> struct all_view_fn { - template <typename _T, typename sycl::access::mode AccMode = sycl::access::mode::read> + template <typename _T> _ONEDPL_CONSTEXPR_FUN oneapi::dpl::__ranges::all_view<_T, AccMode> operator()(sycl::buffer<_T, 1> __buf, typename ::std::iterator_traits<_T*>::difference_type __offset = 0, typename ::std::iterator_traits<_T*>::difference_type __n = 0) const @@ -91,19 +92,20 @@ struct all_view_fn return oneapi::dpl::__ranges::all_view<_T, AccMode>(__buf, __offset, __n); } - template <typename _T, typename sycl::access::mode AccMode> - _ONEDPL_CONSTEXPR_FUN oneapi::dpl::__ranges::all_view<_T, AccMode> - operator()(oneapi::dpl::__ranges::all_view<_T, AccMode> __view) const + template <typename _R> + auto + operator()(_R&& __r) const -> decltype(::std::forward<_R>(__r)) { - return __view; + return ::std::forward<_R>(__r); } }; namespace views { - _ONEDPL_CONSTEXPR_VAR all_view_fn all; -} +_ONEDPL_CONSTEXPR_VAR all_view_fn<sycl::access::mode::read> all_read; +_ONEDPL_CONSTEXPR_VAR all_view_fn<sycl::access::mode::write> all_write; +} // namespace views //all_view traits diff --git a/include/oneapi/dpl/pstl/ranges_defs.h b/include/oneapi/dpl/pstl/ranges_defs.h index cb57ad104c57..0232a17187b9 100644 --- a/include/oneapi/dpl/pstl/ranges_defs.h +++ b/include/oneapi/dpl/pstl/ranges_defs.h @@ -43,6 +43,8 @@ using nano::ranges::transform_view; namespace views { using oneapi::dpl::__ranges::views::all; +using oneapi::dpl::__ranges::views::all_read; +using oneapi::dpl::__ranges::views::all_write; using nano::views::drop; using nano::views::iota; diff --git a/test/parallel_api/ranges/copy_ranges_factory_sycl.pass.cpp b/test/parallel_api/ranges/copy_ranges_factory_sycl.pass.cpp index 5bbbfeae6996..35af0191ffa8 100644 --- a/test/parallel_api/ranges/copy_ranges_factory_sycl.pass.cpp +++ b/test/parallel_api/ranges/copy_ranges_factory_sycl.pass.cpp @@ -39,11 +39,14 @@ main() { sycl::buffer<int> B(data2, sycl::range<1>(max_n)); + sycl::buffer<int> C(max_n); auto view = iota_view(0, max_n) | views::transform(lambda1); + copy(TestUtils::default_dpcpp_policy, view, C); //check passing a buffer for writting + auto range_res = all_view<int, sycl::access::mode::write>(B); - copy(TestUtils::default_dpcpp_policy, view, range_res); + copy(TestUtils::default_dpcpp_policy, C, range_res); //check passing a buffer for reading } //check result diff --git a/test/parallel_api/ranges/copy_ranges_sycl.pass.cpp b/test/parallel_api/ranges/copy_ranges_sycl.pass.cpp index 4c66eea26286..bdf2fb6ad67d 100644 --- a/test/parallel_api/ranges/copy_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/copy_ranges_sycl.pass.cpp @@ -40,13 +40,15 @@ main() { sycl::buffer<int> A(data, sycl::range<1>(max_n)); sycl::buffer<int> B(data2, sycl::range<1>(max_n)); + sycl::buffer<int> C(max_n); auto sv = all_view(A); auto view = views::reverse(sv) | views::transform(lambda1); auto range_res = all_view<int, sycl::access::mode::write>(B); - copy(TestUtils::default_dpcpp_policy, view, range_res); + copy(TestUtils::default_dpcpp_policy, view, C); //check passing a buffer for writting + copy(TestUtils::default_dpcpp_policy, C, range_res); //check passing a buffer for reading } //check result diff --git a/test/parallel_api/ranges/exclusive_scan_ranges_sycl.pass.cpp b/test/parallel_api/ranges/exclusive_scan_ranges_sycl.pass.cpp index 73562056b178..bfefc9a1d3ef 100644 --- a/test/parallel_api/ranges/exclusive_scan_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/exclusive_scan_ranges_sycl.pass.cpp @@ -45,13 +45,12 @@ main() auto view = ranges::all_view<int, sycl::access::mode::read>(A); auto view_res1 = ranges::all_view<int, sycl::access::mode::write>(B1); - auto view_res2 = ranges::all_view<int, sycl::access::mode::write>(B2); auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - ranges::exclusive_scan(exec, view, view_res1, 100); - ranges::exclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, view_res2, 100, ::std::plus<int>()); + ranges::exclusive_scan(exec, A, view_res1, 100); + ranges::exclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, B2, 100, ::std::plus<int>()); } //check result diff --git a/test/parallel_api/ranges/find_end_ranges_sycl.pass.cpp b/test/parallel_api/ranges/find_end_ranges_sycl.pass.cpp index 55135e8d58c9..18520e8c19c9 100644 --- a/test/parallel_api/ranges/find_end_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/find_end_ranges_sycl.pass.cpp @@ -46,6 +46,7 @@ main() auto view_a = all_view(A); auto view_b = all_view(B); res = find_end(TestUtils::default_dpcpp_policy, view_a, view_b); + res = find_end(TestUtils::default_dpcpp_policy, A, B); //check passing sycl buffer directly } //check result diff --git a/test/parallel_api/ranges/find_first_of_ranges_sycl.pass.cpp b/test/parallel_api/ranges/find_first_of_ranges_sycl.pass.cpp index a905eff95f34..ffec8809ba6f 100644 --- a/test/parallel_api/ranges/find_first_of_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/find_first_of_ranges_sycl.pass.cpp @@ -46,6 +46,7 @@ main() auto view_a = all_view(A); auto view_b = all_view(B); res = find_first_of(TestUtils::default_dpcpp_policy, view_a, view_b); + res = find_first_of(TestUtils::default_dpcpp_policy, A, B); //check passing sycl buffer directly } //check result diff --git a/test/parallel_api/ranges/find_ranges_sycl.pass.cpp b/test/parallel_api/ranges/find_ranges_sycl.pass.cpp index 16ea6e92b6ad..e3e41ceb4978 100644 --- a/test/parallel_api/ranges/find_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/find_ranges_sycl.pass.cpp @@ -48,9 +48,12 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - res1 = find(exec, view, val); - res2 = find_if(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, [val](auto a) { return a == val;}); - res3 = find_if_not(make_new_policy<new_kernel_name<Policy, 1>>(exec), view, [val](auto a) { return a >= 0;}); + res1 = find(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, val); //check passing all_view + res1 = find(make_new_policy<new_kernel_name<Policy, 1>>(exec), A, val); //check passing sycl::buffer directly + res2 = find_if(make_new_policy<new_kernel_name<Policy, 2>>(exec), view, [val](auto a) { return a == val;}); + res2 = find_if(make_new_policy<new_kernel_name<Policy, 3>>(exec), A, [val](auto a) { return a == val;}); + res3 = find_if_not(make_new_policy<new_kernel_name<Policy, 4>>(exec), view, [val](auto a) { return a >= 0;}); + res3 = find_if_not(make_new_policy<new_kernel_name<Policy, 5>>(exec), A, [val](auto a) { return a >= 0;}); } //check result diff --git a/test/parallel_api/ranges/for_each_ranges_sycl.pass.cpp b/test/parallel_api/ranges/for_each_ranges_sycl.pass.cpp index ea8eb2300f5b..1d67c33f0bfc 100644 --- a/test/parallel_api/ranges/for_each_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/for_each_ranges_sycl.pass.cpp @@ -31,6 +31,7 @@ main() #if _ENABLE_RANGES_TESTING constexpr int max_n = 10; int data[max_n] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; + int expected[max_n] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; auto lambda1 = [](auto& val) { return val = val * val; }; @@ -38,14 +39,13 @@ main() { sycl::buffer<int> A(data, sycl::range<1>(max_n)); - - auto view = all_view<int, sycl::access::mode::read_write>(A); - for_each(TestUtils::default_dpcpp_policy, view, lambda1); + for_each(TestUtils::default_dpcpp_policy, all_view<int, sycl::access::mode::read_write>(A), lambda1); + for_each(TestUtils::default_dpcpp_policy, A, lambda1); //check with passing sycl::buffer directly } //check result - int expected[max_n]; - ::std::transform(data, data + max_n, expected, lambda1); + ::std::for_each(expected, expected + max_n, lambda1); + ::std::for_each(expected, expected + max_n, lambda1); EXPECT_EQ_N(expected, data, max_n, "wrong effect from for_each with sycl ranges"); #endif //_ENABLE_RANGES_TESTING diff --git a/test/parallel_api/ranges/inclusive_scan_ranges_sycl.pass.cpp b/test/parallel_api/ranges/inclusive_scan_ranges_sycl.pass.cpp index 110e8c8db6fc..b6d2c3cda606 100644 --- a/test/parallel_api/ranges/inclusive_scan_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/inclusive_scan_ranges_sycl.pass.cpp @@ -52,8 +52,8 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - ranges::inclusive_scan(exec, view, view_res1); - ranges::inclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, view_res2, ::std::plus<int>()); + ranges::inclusive_scan(exec, A, view_res1); + ranges::inclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, B2, ::std::plus<int>()); ranges::inclusive_scan(make_new_policy<new_kernel_name<Policy, 1>>(exec), view, view_res3, ::std::plus<int>(), 100); } diff --git a/test/parallel_api/ranges/is_sorted_ranges_sycl.pass.cpp b/test/parallel_api/ranges/is_sorted_ranges_sycl.pass.cpp index ef1a0e490e33..77d6a0a71e8d 100644 --- a/test/parallel_api/ranges/is_sorted_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/is_sorted_ranges_sycl.pass.cpp @@ -45,7 +45,7 @@ main() using Policy = decltype(TestUtils::default_dpcpp_policy); res1 = is_sorted(exec, all_view(A)); - res2 = is_sorted(make_new_policy<new_kernel_name<Policy, 0>>(exec), all_view(B)); + res2 = is_sorted(make_new_policy<new_kernel_name<Policy, 0>>(exec), B); } //check result diff --git a/test/parallel_api/ranges/is_sorted_until_ranges_sycl.pass.cpp b/test/parallel_api/ranges/is_sorted_until_ranges_sycl.pass.cpp index f42f0c47fec5..ddb15a970d4d 100644 --- a/test/parallel_api/ranges/is_sorted_until_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/is_sorted_until_ranges_sycl.pass.cpp @@ -47,7 +47,7 @@ main() using Policy = decltype(TestUtils::default_dpcpp_policy); res1 = is_sorted_until(exec, view); - res2 = is_sorted_until(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, [](auto a, auto b) { return a < b; }); + res2 = is_sorted_until(make_new_policy<new_kernel_name<Policy, 0>>(exec), A, [](auto a, auto b) { return a < b; }); } //check result diff --git a/test/parallel_api/ranges/merge_ranges_sycl.pass.cpp b/test/parallel_api/ranges/merge_ranges_sycl.pass.cpp index 76fe8c246cc1..ce05bcfda19a 100644 --- a/test/parallel_api/ranges/merge_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/merge_ranges_sycl.pass.cpp @@ -50,7 +50,7 @@ main() auto exec = TestUtils::default_dpcpp_policy; merge(exec, all_view(A), all_view(B), all_view<T, sycl::access::mode::write>(D)); - merge(oneapi::dpl::execution::make_device_policy<class merge_2>(exec), all_view(A), all_view(B), all_view<T, sycl::access::mode::write>(E), ::std::less<T>()); + merge(oneapi::dpl::execution::make_device_policy<class merge_2>(exec), A, B, E, ::std::less<T>()); //check passing sycl buffers directly } //check result diff --git a/test/parallel_api/ranges/minmax_ranges_sycl.pass.cpp b/test/parallel_api/ranges/minmax_ranges_sycl.pass.cpp index 90f490feb0ae..1cea7e642ca8 100644 --- a/test/parallel_api/ranges/minmax_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/minmax_ranges_sycl.pass.cpp @@ -51,15 +51,15 @@ main() using Policy = decltype(TestUtils::default_dpcpp_policy); //min element - res1 = min_element(exec, view); + res1 = min_element(exec, A); res2 = min_element(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, ::std::less<int>()); res3 = min_element(make_new_policy<new_kernel_name<Policy, 1>>(exec), view | views::take(1)); //max_element - res4 = max_element(make_new_policy<new_kernel_name<Policy, 2>>(exec), view); + res4 = max_element(make_new_policy<new_kernel_name<Policy, 2>>(exec), A); res5 = max_element(make_new_policy<new_kernel_name<Policy, 3>>(exec), view, ::std::less<int>()); - res_minmax1 = minmax_element(make_new_policy<new_kernel_name<Policy, 4>>(exec), view); + res_minmax1 = minmax_element(make_new_policy<new_kernel_name<Policy, 4>>(exec), A); res_minmax2 = minmax_element(make_new_policy<new_kernel_name<Policy, 5>>(exec), view, ::std::less<int>()); } diff --git a/test/parallel_api/ranges/reduce_ranges_sycl.pass.cpp b/test/parallel_api/ranges/reduce_ranges_sycl.pass.cpp index 5a21cebecdad..2c12f14773d1 100644 --- a/test/parallel_api/ranges/reduce_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/reduce_ranges_sycl.pass.cpp @@ -43,7 +43,7 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - res1 = reduce(exec, view); + res1 = reduce(exec, A); res2 = reduce(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, 100); res3 = reduce(make_new_policy<new_kernel_name<Policy, 1>>(exec), view, 100, ::std::plus<int>()); } diff --git a/test/parallel_api/ranges/search_n_ranges_sycl.pass.cpp b/test/parallel_api/ranges/search_n_ranges_sycl.pass.cpp index 97e89474439c..c8f01d4e4e67 100644 --- a/test/parallel_api/ranges/search_n_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/search_n_ranges_sycl.pass.cpp @@ -35,18 +35,20 @@ main() const int n_val = 3; const int idx = 4; const int val = data[idx]; - int res = -1; + int res1 = -1, res2 = -1; using namespace oneapi::dpl::experimental::ranges; { sycl::buffer<int> A(data, sycl::range<1>(count)); auto view_a = all_view(A); - res = search_n(TestUtils::default_dpcpp_policy, view_a, n_val, val, [](auto a, auto b) { return a == b; }); + res1 = search_n(TestUtils::default_dpcpp_policy, view_a, n_val, val, [](auto a, auto b) { return a == b; }); + res2 = search_n(TestUtils::default_dpcpp_policy, A, n_val, val, [](auto a, auto b) { return a == b; }); } //check result - EXPECT_TRUE(res == idx, "wrong effect from 'search_n' with sycl ranges"); + EXPECT_TRUE(res1 == idx, "wrong effect from 'search_n' with sycl ranges"); + EXPECT_TRUE(res2 == idx, "wrong effect from 'search_n' with sycl buffer"); #endif //_ENABLE_RANGES_TESTING diff --git a/test/parallel_api/ranges/search_ranges_sycl.pass.cpp b/test/parallel_api/ranges/search_ranges_sycl.pass.cpp index 755c6f53850a..95bde4590cca 100644 --- a/test/parallel_api/ranges/search_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/search_ranges_sycl.pass.cpp @@ -50,8 +50,8 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - res1 = search(exec, view_a, view_b); - res2 = search(make_new_policy<new_kernel_name<Policy, 0>>(exec), view_a, view_b, [](auto a, auto b) { return a == b; }); + res1 = search(exec, A, view_b); + res2 = search(make_new_policy<new_kernel_name<Policy, 0>>(exec), view_a, B, [](auto a, auto b) { return a == b; }); } //check result diff --git a/test/parallel_api/ranges/sort_ranges_sycl.pass.cpp b/test/parallel_api/ranges/sort_ranges_sycl.pass.cpp index 54f0484f9877..8bffd87db783 100644 --- a/test/parallel_api/ranges/sort_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/sort_ranges_sycl.pass.cpp @@ -42,7 +42,7 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - sort(exec, all_view<int, sycl::access::mode::read_write>(A)); + sort(exec, A); //check passing sycl buffer directly sort(make_new_policy<new_kernel_name<Policy, 0>>(exec), all_view<int, sycl::access::mode::read_write>(B), ::std::greater<int>()); } diff --git a/test/parallel_api/ranges/stable_sort_ranges_sycl.pass.cpp b/test/parallel_api/ranges/stable_sort_ranges_sycl.pass.cpp index 248823470e70..5687980a59b8 100644 --- a/test/parallel_api/ranges/stable_sort_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/stable_sort_ranges_sycl.pass.cpp @@ -42,7 +42,7 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - stable_sort(exec, all_view<int, sycl::access::mode::read_write>(A)); + stable_sort(exec, A); //check passing sycl buffer directly stable_sort(make_new_policy<new_kernel_name<Policy, 0>>(exec), all_view<int, sycl::access::mode::read_write>(B), ::std::greater<int>()); } diff --git a/test/parallel_api/ranges/transform2_ranges_factory_sycl.pass.cpp b/test/parallel_api/ranges/transform2_ranges_factory_sycl.pass.cpp index 2556c047da8f..a908b0a1fe16 100644 --- a/test/parallel_api/ranges/transform2_ranges_factory_sycl.pass.cpp +++ b/test/parallel_api/ranges/transform2_ranges_factory_sycl.pass.cpp @@ -32,6 +32,7 @@ main() constexpr int max_n = 10; int data[max_n] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; int data2[max_n]; + int data3[max_n]; auto lambda1 = [](auto i) { return i * i; }; auto lambda2 = [](auto i, auto j) { return i + j; }; @@ -40,11 +41,13 @@ main() { sycl::buffer<int> B(data2, sycl::range<1>(max_n)); + sycl::buffer<int> C(data3, sycl::range<1>(max_n)); auto view = iota_view(0, max_n) | views::transform(lambda1); auto range_res = all_view<int, sycl::access::mode::write>(B); transform(TestUtils::default_dpcpp_policy, view, view, range_res, lambda2); + transform(TestUtils::default_dpcpp_policy, view, view, C, lambda2); //check passing sycl buffer } //check result @@ -53,6 +56,7 @@ main() ::std::transform(expected, expected + max_n, expected, expected, lambda2); EXPECT_EQ_N(expected, data2, max_n, "wrong effect from trasnform2 with sycl ranges"); + EXPECT_EQ_N(expected, data3, max_n, "wrong effect from trasnform2 with sycl buffer"); #endif //_ENABLE_RANGES_TESTING ::std::cout << TestUtils::done() << ::std::endl; return 0; diff --git a/test/parallel_api/ranges/transform_exclusive_scan_ranges_sycl.pass.cpp b/test/parallel_api/ranges/transform_exclusive_scan_ranges_sycl.pass.cpp index a171b33cb094..c771ee7b3ba2 100644 --- a/test/parallel_api/ranges/transform_exclusive_scan_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/transform_exclusive_scan_ranges_sycl.pass.cpp @@ -34,11 +34,13 @@ main() constexpr int max_n = 10; int data[max_n] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; int data1[max_n]; + int data2[max_n]; auto lambda = [](auto i) { return i * i; }; { sycl::buffer<int> A(data, sycl::range<1>(max_n)); sycl::buffer<int> B(data1, sycl::range<1>(max_n)); + sycl::buffer<int> C(data2, sycl::range<1>(max_n)); using namespace oneapi::dpl::experimental; @@ -46,6 +48,7 @@ main() auto view_res = ranges::all_view<int, sycl::access::mode::write>(B); ranges::transform_exclusive_scan(TestUtils::default_dpcpp_policy, view, view_res, 100, ::std::plus<int>(), lambda); + ranges::transform_exclusive_scan(TestUtils::default_dpcpp_policy, A, C, 100, ::std::plus<int>(), lambda); } //check result @@ -53,6 +56,7 @@ main() ::std::transform_exclusive_scan(oneapi::dpl::execution::seq, data, data + max_n, expected, 100, ::std::plus<int>(), lambda); EXPECT_EQ_N(expected, data1, max_n, "wrong effect from transform_exclusive_scan with init, sycl ranges"); + EXPECT_EQ_N(expected, data2, max_n, "wrong effect from transform_exclusive_scan with init, sycl buffers"); #endif //_ENABLE_RANGES_TESTING ::std::cout << TestUtils::done() << ::std::endl; diff --git a/test/parallel_api/ranges/transform_inclusive_scan_ranges_sycl.pass.cpp b/test/parallel_api/ranges/transform_inclusive_scan_ranges_sycl.pass.cpp index b28a271a8bff..01877a2d0989 100644 --- a/test/parallel_api/ranges/transform_inclusive_scan_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/transform_inclusive_scan_ranges_sycl.pass.cpp @@ -47,13 +47,12 @@ main() auto view = ranges::all_view<int, sycl::access::mode::read>(A); auto view_res1 = ranges::all_view<int, sycl::access::mode::write>(B1); - auto view_res2 = ranges::all_view<int, sycl::access::mode::write>(B2); auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - ranges::transform_inclusive_scan(exec, view, view_res1, ::std::plus<int>(), lambda); - ranges::transform_inclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, view_res2, ::std::plus<int>(), lambda, init); + ranges::transform_inclusive_scan(exec, A, view_res1, ::std::plus<int>(), lambda); + ranges::transform_inclusive_scan(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, B2, ::std::plus<int>(), lambda, init); } //check result diff --git a/test/parallel_api/ranges/transform_ranges_factory_sycl.pass.cpp b/test/parallel_api/ranges/transform_ranges_factory_sycl.pass.cpp index 7a4133499469..e16aa03e93e8 100644 --- a/test/parallel_api/ranges/transform_ranges_factory_sycl.pass.cpp +++ b/test/parallel_api/ranges/transform_ranges_factory_sycl.pass.cpp @@ -32,6 +32,7 @@ main() constexpr int max_n = 10; int data[max_n] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9}; int data2[max_n]; + int data3[max_n]; auto lambda1 = [](auto i) { return i * i; }; auto lambda2 = [](auto i) { return i + 200; }; @@ -40,11 +41,13 @@ main() { sycl::buffer<int> B(data2, sycl::range<1>(max_n)); + sycl::buffer<int> C(data3, sycl::range<1>(max_n)); auto view = iota_view(0, max_n) | views::transform(lambda1); auto range_res = all_view<int, sycl::access::mode::write>(B); transform(TestUtils::default_dpcpp_policy, view, range_res, lambda2); + transform(TestUtils::default_dpcpp_policy, view, C, lambda2); //check passing sycl buffer } //check result @@ -53,6 +56,7 @@ main() ::std::transform(expected, expected + max_n, expected, lambda2); EXPECT_EQ_N(expected, data2, max_n, "wrong effect from trasnform with sycl ranges"); + EXPECT_EQ_N(expected, data3, max_n, "wrong effect from trasnform with sycl buffer"); #endif //_ENABLE_RANGES_TESTING ::std::cout << TestUtils::done() << ::std::endl; return 0; diff --git a/test/parallel_api/ranges/transform_reduce_ranges_sycl.pass.cpp b/test/parallel_api/ranges/transform_reduce_ranges_sycl.pass.cpp index ccda96157899..db66500dd858 100644 --- a/test/parallel_api/ranges/transform_reduce_ranges_sycl.pass.cpp +++ b/test/parallel_api/ranges/transform_reduce_ranges_sycl.pass.cpp @@ -45,8 +45,8 @@ main() auto exec = TestUtils::default_dpcpp_policy; using Policy = decltype(TestUtils::default_dpcpp_policy); - res1 = oneapi::dpl::experimental::ranges::transform_reduce(exec, view, view, 0); - res2 = oneapi::dpl::experimental::ranges::transform_reduce(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, view, 0, ::std::plus<int>(), ::std::multiplies<int>()); + res1 = oneapi::dpl::experimental::ranges::transform_reduce(exec, A, view, 0); + res2 = oneapi::dpl::experimental::ranges::transform_reduce(make_new_policy<new_kernel_name<Policy, 0>>(exec), view, A, 0, ::std::plus<int>(), ::std::multiplies<int>()); res3 = oneapi::dpl::experimental::ranges::transform_reduce(make_new_policy<new_kernel_name<Policy, 1>>(exec), view, 0, ::std::plus<int>(), lambda1); } |