[Draft] Scan changes to improve numerical type requirements

include/oneapi/dpl/pstl/glue_numeric_impl.h

@@ -40,16 +40,16 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, _Tp>
 reduce(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIterator __last, _Tp __init,
        _BinaryOperation __binary_op)
 {
-    return transform_reduce(std::forward<_ExecutionPolicy>(__exec), __first, __last, __init, __binary_op,
+    return transform_reduce(std::forward<_ExecutionPolicy>(__exec), __first, __last, std::move(__init), __binary_op,
                             oneapi::dpl::identity{});
 }
 
 template <class _ExecutionPolicy, class _ForwardIterator, class _Tp>
 oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, _Tp>
 reduce(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIterator __last, _Tp __init)
 {
-    return transform_reduce(std::forward<_ExecutionPolicy>(__exec), __first, __last, __init, std::plus<_Tp>(),
-                            oneapi::dpl::identity{});
+    return transform_reduce(std::forward<_ExecutionPolicy>(__exec), __first, __last, std::move(__init),
+                            std::plus<_Tp>(), oneapi::dpl::identity{});
 }
 
 template <class _ExecutionPolicy, class _ForwardIterator>
@@ -74,7 +74,7 @@ transform_reduce(_ExecutionPolicy&& __exec, _ForwardIterator1 __first1, _Forward
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first1, __first2);
 
     return oneapi::dpl::__internal::__pattern_transform_reduce(
-        __dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first1, __last1, __first2, __init,
+        __dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec), __first1, __last1, __first2, std::move(__init),
         ::std::plus<_InputType>(), ::std::multiplies<_InputType>());
 }
 
@@ -87,8 +87,8 @@ transform_reduce(_ExecutionPolicy&& __exec, _ForwardIterator1 __first1, _Forward
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first1, __first2);
 
     return oneapi::dpl::__internal::__pattern_transform_reduce(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec),
-                                                               __first1, __last1, __first2, __init, __binary_op1,
-                                                               __binary_op2);
+                                                               __first1, __last1, __first2, std::move(__init),
+                                                               __binary_op1, __binary_op2);
 }
 
 template <class _ExecutionPolicy, class _ForwardIterator, class _Tp, class _BinaryOperation, class _UnaryOperation>
@@ -99,7 +99,8 @@ transform_reduce(_ExecutionPolicy&& __exec, _ForwardIterator __first, _ForwardIt
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first);
 
     return oneapi::dpl::__internal::__pattern_transform_reduce(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec),
-                                                               __first, __last, __init, __binary_op, __unary_op);
+                                                               __first, __last, std::move(__init), __binary_op,
+                                                               __unary_op);
 }
 
 // [exclusive.scan]
@@ -109,8 +110,8 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, _Forward
 exclusive_scan(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _ForwardIterator1 __last,
                _ForwardIterator2 __result, _Tp __init)
 {
-    return transform_exclusive_scan(std::forward<_ExecutionPolicy>(__exec), __first, __last, __result, __init,
-                                    std::plus<_Tp>(), oneapi::dpl::identity{});
+    return transform_exclusive_scan(std::forward<_ExecutionPolicy>(__exec), __first, __last, __result,
+                                    std::move(__init), std::plus<_Tp>(), oneapi::dpl::identity{});
 }
 
 #if !_ONEDPL_EXCLUSIVE_SCAN_WITH_BINARY_OP_AMBIGUITY
@@ -119,40 +120,44 @@ oneapi::dpl::__internal::__enable_if_execution_policy<_ExecutionPolicy, _Forward
 exclusive_scan(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _ForwardIterator1 __last,
                _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(std::forward<_ExecutionPolicy>(__exec), __first, __last, __result, __init,
-                                    __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(std::forward<_ExecutionPolicy>(__exec), __first, __last, __result,
+                                    std::move(__init), __binary_op, oneapi::dpl::identity{});
 }
 #else
 template <class _ForwardIterator1, class _ForwardIterator2, class _Tp, class _BinaryOperation>
 _ForwardIterator2
 exclusive_scan(oneapi::dpl::execution::sequenced_policy __exec, _ForwardIterator1 __first, _ForwardIterator1 __last,
                _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 
 template <class _ForwardIterator1, class _ForwardIterator2, class _Tp, class _BinaryOperation>
 _ForwardIterator2
 exclusive_scan(oneapi::dpl::execution::unsequenced_policy __exec, _ForwardIterator1 __first, _ForwardIterator1 __last,
                _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 
 template <class _ForwardIterator1, class _ForwardIterator2, class _Tp, class _BinaryOperation>
 _ForwardIterator2
 exclusive_scan(oneapi::dpl::execution::parallel_policy __exec, _ForwardIterator1 __first, _ForwardIterator1 __last,
                _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 
 template <class _ForwardIterator1, class _ForwardIterator2, class _Tp, class _BinaryOperation>
 _ForwardIterator2
 exclusive_scan(oneapi::dpl::execution::parallel_unsequenced_policy __exec, _ForwardIterator1 __first,
                _ForwardIterator1 __last, _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 
 #    if _ONEDPL_BACKEND_SYCL
@@ -161,7 +166,8 @@ _ForwardIterator2
 exclusive_scan(const oneapi::dpl::execution::device_policy<PolicyParams...>& __exec, _ForwardIterator1 __first,
                _ForwardIterator1 __last, _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 
 #        if _ONEDPL_FPGA_DEVICE
@@ -171,7 +177,8 @@ _ForwardIterator2
 exclusive_scan(const oneapi::dpl::execution::fpga_policy<factor, KernelName>& __exec, _ForwardIterator1 __first,
                _ForwardIterator1 __last, _ForwardIterator2 __result, _Tp __init, _BinaryOperation __binary_op)
 {
-    return transform_exclusive_scan(__exec, __first, __last, __result, __init, __binary_op, oneapi::dpl::identity{});
+    return transform_exclusive_scan(__exec, __first, __last, __result, std::move(__init), __binary_op,
+                                    oneapi::dpl::identity{});
 }
 #        endif // _ONEDPL_FPGA_DEVICE
 #    endif     // _ONEDPL_BACKEND_SYCL
@@ -205,7 +212,7 @@ inclusive_scan(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _ForwardIte
                _ForwardIterator2 __result, _BinaryOperation __binary_op, _Tp __init)
 {
     return transform_inclusive_scan(std::forward<_ExecutionPolicy>(__exec), __first, __last, __result, __binary_op,
-                                    oneapi::dpl::identity{}, __init);
+                                    oneapi::dpl::identity{}, std::move(__init));
 }
 
 // [transform.exclusive.scan]
@@ -220,8 +227,8 @@ transform_exclusive_scan(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first, __result);
 
     return oneapi::dpl::__internal::__pattern_transform_scan(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec),
-                                                             __first, __last, __result, __unary_op, __init, __binary_op,
-                                                             /*inclusive=*/::std::false_type());
+                                                             __first, __last, __result, __unary_op, std::move(__init),
+                                                             __binary_op, /*inclusive=*/::std::false_type());
 }
 
 // [transform.inclusive.scan]
@@ -236,8 +243,8 @@ transform_inclusive_scan(_ExecutionPolicy&& __exec, _ForwardIterator1 __first, _
     const auto __dispatch_tag = oneapi::dpl::__internal::__select_backend(__exec, __first, __result);
 
     return oneapi::dpl::__internal::__pattern_transform_scan(__dispatch_tag, ::std::forward<_ExecutionPolicy>(__exec),
-                                                             __first, __last, __result, __unary_op, __init, __binary_op,
-                                                             /*inclusive=*/::std::true_type());
+                                                             __first, __last, __result, __unary_op, std::move(__init),
+                                                             __binary_op, /*inclusive=*/::std::true_type());
 }
 
 template <class _ExecutionPolicy, class _ForwardIterator1, class _ForwardIterator2, class _UnaryOperation,

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h

@@ -667,9 +667,15 @@ struct __result_and_scratch_storage : __result_and_scratch_storage_base
         }
         else if (__supports_USM_device)
         {
-            _T __tmp;
-            __q.memcpy(&__tmp, __scratch_buf.get() + __scratch_n + _Idx, 1 * sizeof(_T)).wait();
-            return __tmp;
+            static_assert(sycl::is_device_copyable_v<_T>,
+                          "The type _T must be device copyable to use __result_and_scratch_storage.");
+            // Avoid default constructor for _T, we know that _T is device copyable and therefore a copy construction
+            // is equivalent to a bitwise copy
+            _T* __tmp = static_cast<_T*>(malloc(sizeof(_T)));
+            __q.memcpy(__tmp, __scratch_buf.get() + __scratch_n + _Idx, 1 * sizeof(_T)).wait();
+            _T __return_tmp = std::move(*__tmp);
+            free(__tmp);
+            return __return_tmp;
         }
         else
         {