PR NVIDIA/cub#218 fixes this CUB's radix sort. We should:

• Check whether Thrust's other backends handle this case correctly.
• Provide a guarantee of this in the stable_sort documentation.
• Add regression tests to enforce this on all backends.

Report needed documentation

We do not have documentation specifying the different treelite Operator values that FIL supports. (https://github.com/dmlc/treelite/blob/46c8390aed4491ea97a017d447f921efef9f03ef/include/treelite/base.h#L40)

Report needed documentation
https://github.com/rapidsai/cuml/blob/branch-0.15/cpp/test/sg/fil_test.cu
There are multiple places in the fil_test.cu file

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,
        OffsetIterato