NVIDIA / NVIDIA/cccl

Specialize algorithms to treat "contiguous" iterators like raw pointers

Open
#914 1 comment 1 reaction 1 assignee Claimed by @ericniebler View on GitHub
cub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

Most CUB algorithms accept generic Iterator ranges as inputs.

When the input is a raw pointer, CUB may apply additional optimizations to exploit the fact that the elements in the range are contiguous in memory, e.g., using vectorized loads.

These same optimizations cannot be made on a general iterator as there are no guarantees about the elements being contiguous in memory (or even existing in memory).

However, many iterators do refer to elements that are contiguous in memory and could benefit from the same optimizations applied to raw pointers.

We should detect those iterators when we can and treat them as if there were pointers.

Thrust already has logic to detect "contiguous" iterators, but there is no magic. We can't detect all possible contiguous iterators. It's basically just a bunch of specializations for common iterator types that we know are contiguous, e.g., `thrust::device_vector::iterator`, `std::vector::iterator`.

If you have an iterator type `It` that you know is contiguous that Thrust doesn't already know about, you can "proclaim" it as contiguous via `THRUST_PROCLAIM_CONTIGUOUS_ITERATOR(It)`

See: https://github.com/NVIDIA/thrust/blob/65fbe23ab95d58966a2bc44245c084576f093b71/thrust/type_traits/is_contiguous_iterator.h#L68-L216

See also [`std::contiguous_iterator`](https://en.cppreference.com/w/cpp/iterator/contiguous_iterator) concept added in C++20.

We should exploit this same logic in CUB.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.