Allow segmented problems to have different types for offset iterator #229

devavret · 2020-10-29T10:34:39Z

Segmented reduce uses the same template type OffsetIteratorT for begin and end offsets

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceSegmentedReduce::Sum
    (   void *              d_temp_storage,
        size_t &            temp_storage_bytes,
        InputIteratorT      d_in,
        OutputIteratorT     d_out,
        int                 num_segments,
        OffsetIteratorT     d_begin_offsets,  ///<<<
        OffsetIteratorT     d_end_offsets,    ///<<<
        cudaStream_t        stream = 0,
        bool                debug_synchronous = false 
    )

This precludes use of thrust transform iterators to generate the offsets. The implementation suggests that the iterators are only dereferenced to get the begin and end iterators independently per-block and as long as the dereferenced value is the same type, it should be ok for the iterators to be a different type.

Check the attached src for and example use case.
cub_segment_repro.cu.zip

allisonvacanti added enhancement good first issue labels Oct 30, 2020

allisonvacanti added this to the 1.12.0 milestone Oct 30, 2020

Nov	DEC	Jan
	01
2019	2020	2021

NVIDIA / cub

Allow segmented problems to have different types for offset iterator #229

Allow segmented problems to have different types for offset iterator #229

devavret commented Oct 29, 2020

NVIDIA / cub

Join GitHub today

GitHub is where the world builds software

Allow segmented problems to have different types for offset iterator #229

Allow segmented problems to have different types for offset iterator #229

Comments

devavret commented Oct 29, 2020

Essential cookies

Always active

Analytics cookies