Getting starting index of groups within dataset - ReduceByKey

[stndbye]

Hello,

I'm just learning about CUB, so i don't have much knowledge about the library.

I have a list of non unique sorted keys, and i would like to produce the first index for each key.
ReduceByKey sounded perfect for the job - created an additional array that is passed on as d_values_in, and contains [0,1,2...n].
Using a min op i have what i need.
It just feels so wrong having to store a linear counter in my precious memory and then accessing global memory for it.

Is there a way to avoid storing the value array - and somehow just feed the index of the current key as an input value? Or is there another method better suited for this task?

[elstehle]

Hello @stndbye and thanks for reaching out!

Great to hear you are getting started with CUB!

You have already correctly identified a great solution to the problem you're trying to solve by using cub's cub::DeviceReduce::ReduceByKey algorithm.

And you have correctly questioned whether you would really need to materialize the integer sequence [0, 1, 2, ..., n]. Good news: You don't have to materialize that sequence.

CUB offers fancy iterators to accomplish tasks like the one you have outlined. In your specific case, you can make use of cub::CountingInputIterator that returns i when the i-th element is dereferenced.

I have put together a sample usage here.

Worth noting, most CUB algorithms support fancy iterators, and you can use fancy iterators wherever you would use a CUDA array.

I hope this helps. Please let us know if you have any follow up questions.

[pauleonix]

With there being two redundant implementations of many fancy iterators in CCCL right now, is there any good reason of using the CUB versions over the Thrust versions with their nice make_* factory functions? Will they be superseded by a e.g. cuda::std::views::iota in libcu++?

[pauleonix]

Or maybe cuda::views::iota which can only create a common range/view as sentinels would not work too well with parallel/GPU computing.

[miscco]

We are currently in the process of implementing std::ranges, so there will be a std::views::iota soon ^TM

[elstehle]

I am not aware of any reason to prefer CUB's fancy iterators over thrust's fancy iterators. So, for completeness, @stndbye, note that you can just as well use thrust's fancy iterators in places where you can use CUB's fancy iterators

[stndbye]

Thanks for your help, works like a charm!

[elstehle]

Thanks for confirming the solution worked for you. Closing the issue as resolved.