Address comments in reduce-then-scan based implementation

Signed-off-by: Matthew Michel <[email protected]>

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

@@ -802,6 +802,9 @@ struct __gen_red_by_seg_reduce_input
         const auto __in_keys = std::get<0>(__in_rng.tuple());
         const auto __in_vals = std::get<1>(__in_rng.tuple());
         using _ValueType = oneapi::dpl::__internal::__value_t<decltype(__in_vals)>;
+        // The first segment start (index 0) is not marked with a 1. This is because we need the first
+        // segment's key and value output index to be 0. We begin marking new segments only after the
+        // first.
         const std::size_t __new_seg_mask = __id > 0 && !__binary_pred(__in_keys[__id - 1], __in_keys[__id]);
         return oneapi::dpl::__internal::make_tuple(__new_seg_mask, _ValueType{__in_vals[__id]});
     }
@@ -845,7 +848,7 @@ struct __gen_red_by_seg_scan_input
                 oneapi::dpl::__internal::make_tuple(__new_seg_mask, _ValueType{__in_vals[__id]}), true, __current_key,
                 __current_key); // Passing __current_key as the next key for the last element is a placeholder
         }
-        else
+        else // __id == 0
         {
             const _KeyType& __next_key = __in_keys[__id + 1];
             return oneapi::dpl::__internal::make_tuple(
@@ -867,7 +870,7 @@ struct __red_by_seg_op
     {
         using std::get;
         // The left-hand side has processed elements from the same segment, so update the reduction value.
-        if (std::get<0>(__rhs_tup) == 0)
+        if (get<0>(__rhs_tup) == 0)
         {
             return oneapi::dpl::__internal::make_tuple(get<0>(__lhs_tup),
                                                        __binary_op(get<1>(__lhs_tup), get<1>(__rhs_tup)));
@@ -897,6 +900,11 @@ struct __write_red_by_seg
         const bool __is_seg_end = get<1>(__tup);
         const std::size_t __out_idx = get<0>(get<0>(__tup));
 
+        // With the exception of the first key which is output by index 0, the first key in each segment is written
+        // by the work item that outputs the previous segment's reduction value. This is because the reduce_by_segment
+        // API requires that the first key in a segment is output and is important for when keys in a segment might not
+        // be the same (but satisfy the predicate). The last segment does not output a key as there are no future
+        // segments process.
         if (__id == 0)
             __out_keys[0] = __current_key;
         if (__is_seg_end)

test/general/implementation_details/device_copyable.pass.cpp

@@ -156,9 +156,9 @@ test_device_copyable()
         "__gen_transform_input is not device copyable with device copyable types");
 
     //__gen_red_by_seg_reduce_input
-    static_assert(
-        sycl::is_device_copyable_v<oneapi::dpl::__par_backend_hetero::__gen_red_by_seg_reduce_input<noop_device_copyable>>,
-        "__gen_red_by_seg_reduce_input is not device copyable with device copyable types");
+    static_assert(sycl::is_device_copyable_v<
+                      oneapi::dpl::__par_backend_hetero::__gen_red_by_seg_reduce_input<binary_op_device_copyable>>,
+                  "__gen_red_by_seg_reduce_input is not device copyable with device copyable types");
 
     //__gen_red_by_seg_scan_input
     static_assert(