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);
     }