Offset-based segmented reduce (CSR / CUB DeviceSegmentedReduce equivalent)

[zjin-lcf]

Summary

Please consider adding an offset-based (CSR-style) segmented reduction to oneDPL’s Parallel API.

oneDPL already provides oneapi::dpl::reduce_by_segment, which reduces contiguous runs of equal keys. That is a different abstraction from an offset-based segmented reduce, in which segments are described by begin/end offsets, the number of segments is known a priori, and every segment produces exactly one output, including empty segments.

Motivation

Offset-based segmented reduction is a common building block for:

• Sparse/CSR layouts (offsets[s], offsets[s+1])
• Group-by aggregations after a sort, where group bounds are stored as a prefix-sum / offset array
• Workloads that include empty groups (zero-length segments)

reduce_by_segment cannot express empty groups: a key that never appears yields no output. Representing CSR data as a dense key array also requires extra storage and bandwidth proportional to the number of values, not the number of segments.

CUB (cub::DeviceSegmentedReduce) and hipCUB expose this interface directly. SYCLomatic maps those calls to dpct::device::segmented_reduce, which is a migration helper rather than a oneDPL algorithm, so portable SYCL code still has no library equivalent.

Requested behavior

A device-capable algorithm along the lines of:

reduce_by_offsets(policy,
                  values_first,
                  num_segments,
                  begin_offsets,   // length num_segments
                  end_offsets,     // length num_segments; may alias begin_offsets + 1
                  results,
                  binary_op,
                  identity)

Expected semantics:

• Segment s is the half-open range [begin_offsets[s], end_offsets[s]).
• If end_offsets[s] <= begin_offsets[s], the segment is empty and results[s] = identity.
• Otherwise results[s] is the reduction of that range with binary_op, using identity as the initial value.
• Offset iterators should support 64-bit index types (std::int64_t / std::uint64_t), matching current CUB/hipCUB practice, not only 32-bit int.
• At least the usual operators: plus, minimum, maximum (and a generic associative binary_op if feasible).

A single CSR offset array of length num_segments + 1 should be usable by passing offsets and offsets + 1, as in CUB.

Related APIs

Library	API	Segment definition
oneDPL	reduce_by_segment	Adjacent equal keys
CUB / hipCUB	DeviceSegmentedReduce::{Reduce,Sum,Min,Max}	Begin/end offsets
SYCLomatic / DPCT	dpct::device::segmented_reduce	Begin/end offsets (not oneDPL)

This request is specifically for the offset-based form, not a replacement for reduce_by_segment.

[zjin-lcf]

@danhoeflinger A custom SYCL kernel is available in ORNL/HeCBench#337. Thanks.