[DOC]: Documentation of thrust::reduce_by_key states to select the _first_ key of a range of consecutively equal keys but selects the _last_

[elstehle]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Is this for new documentation, or an update to existing docs?

New

Describe the incorrect/future/missing documentation

Docs of thrust::reduce_by_key incorrectly states to select the first key of a range of consecutively equal keys, however, we select the last key of such range.

reduce_by_key is a generalization of reduce to key-value pairs. For each group of consecutive keys in the range [keys_first, keys_last) that are equal, reduce_by_key copies the first element of the group to the keys_output.

Documentation of cub::DeviceReduce::ReduceByKey makes the same claim. CUB's implementation also looks contrary to this claim.

The following is a reproducer for thrust was provided by Archith Bency:

#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/sequence.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/replace.h>
#include <thrust/functional.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>

#include <cstdint>
#include <iostream>
#include <iomanip>
#include <vector>

template <typename Vector>
void print_vector(const std::string& name, const Vector& v)
{
  typedef typename Vector::value_type T;
  std::cout << name << "  ";
  thrust::copy(v.begin(), v.end(), std::ostream_iterator<T>(std::cout, " "));
  std::cout << std::endl;
}


struct Reduce_pred {
    __host__ __device__ __forceinline__
    bool operator()(const thrust::tuple<uint32_t, uint32_t, float, uint32_t>& valueA,
                const thrust::tuple<uint32_t, uint32_t, float, uint32_t>& valueB) {
    
        
        return (thrust::get<0>(valueA) == thrust::get<0>(valueB)) &&
            (thrust::get<1>(valueA) == thrust::get<1>(valueB));
    } 
};


int main(void)
{

    thrust::device_vector<uint32_t> vecA(std::vector<uint32_t>{0, 0, 0, 0});
    thrust::device_vector<uint32_t> vecB(std::vector<uint32_t>{2, 2, 2, 2});
    thrust::device_vector<float> vecC(std::vector<float>{88.0f, 88.0f, 89.0f, 89.0f});
    
    thrust::device_vector<uint32_t> vecA_reduced(vecA.size());
    thrust::device_vector<uint32_t> vecB_reduced(vecB.size());    
    thrust::device_vector<float> vecC_reduced(vecC.size());
    thrust::device_vector<uint32_t> pickedVecIdx(vecA.size());


    auto inputKeyItBegin = thrust::make_zip_iterator(vecA.begin(), vecB.begin(), vecC.begin(), thrust::make_counting_iterator(uint32_t(0)));
    auto outputKeyItBegin = thrust::make_zip_iterator(vecA_reduced.begin(), vecB_reduced.begin(), vecC_reduced.begin(), pickedVecIdx.begin());


    auto end = thrust::reduce_by_key(inputKeyItBegin,
                                inputKeyItBegin + 4,
                                thrust::make_constant_iterator(uint32_t(0)),
                                outputKeyItBegin,
                                thrust::make_discard_iterator(),
                                Reduce_pred{});
    
    auto num_reduction = end.first - outputKeyItBegin;
    vecA_reduced.resize(num_reduction);
    vecB_reduced.resize(num_reduction);
    vecC_reduced.resize(num_reduction);
    pickedVecIdx.resize(num_reduction);

    print_vector("vecA_reduced", vecA_reduced);
    print_vector("vecB_reduced", vecB_reduced);
    print_vector("vecC_reduced", vecC_reduced);
    print_vector("pickedVecIdx", pickedVecIdx); // prints 3, expected 0 as reduce_by_key should pick the first member of the sole group in this example


    std::cout << "num_reduction : " << num_reduction << std::endl;


    
    return 0;
}

If this is a correction, please provide a link to the incorrect documentation. If this is a new documentation request, please link to where you have looked.

Thrust:
https://thrust.github.io/doc/group__reductions_ga633d78d4cb2650624ec354c9abd0c97f.html#ga633d78d4cb2650624ec354c9abd0c97f

CUB:
https://nvlabs.github.io/cub/structcub_1_1_device_reduce.html#a303ae673ac32825f95912b4bfff8bef1

[kerrybbarnes]

I'm seeing the same issue. Code is

thrust::counting_iterator<int> iter(0);
thrust::constant_iterator<gaia::index_t> one(1);
auto dend = thrust::reduce_by_key(exec,
                                  iter,
                                  iter+num_points, // input key sequence
                                  one,
                                  d_offsets,
                                  d_lengths,
                                  eq_pred).first;
KERNEL_EXEC_SYNCHRONIZE(exec);
size_t num_runs =  dend - d_offsets;
printf("num_runs = %ld\n", num_runs);
print_vector(exec,"d_offsets",d_offsets,num_runs);
print_vector(exec,"d_lengths",d_lengths,num_runs);
printf("----\n");

On both host and device the lengths are correct. For one example the lengths are 3,1,2. The lengths sum the value 1 and are just the number of values in the range of equal values. From them you can infer the offsets should be 0,3,4. But they are 2,3,5. It choosing the last value of iter rather than the first.