[BUG]: `cuda::transform_iterator` silently downgrades `iterator_category` → 2.5–3.3× slowdowns for thrust algorithms that dispatch on it

[ttnghia]

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

Type of Bug

Performance

Component

General CCCL

Describe the bug

Problem

cuda::transform_iterator deliberately sets its classic iterator_category member to input_iterator_tag whenever the wrapped functor returns by value, even when the underlying iterator is random-access. This is spec-compliant for C++20 std::ranges::transform_view::iterator, but it silently breaks downstream Thrust algorithms that dispatch on std::iterator_traits<It>::iterator_category (e.g. thrust::copy → thrust::scatter → thrust::make_permutation_iterator), causing them to fall off their CUB bulk / vectorized fast paths.

In cuDF we observed a 2.5–3.3× GPU-side slowdown after migrating two call sites from thrust::make_transform_iterator to cuda::transform_iterator. One-line revert restored parity with thrust::: rapidsai/cudf#14162.

Observed impact (cuDF scatter regression, RTX 6000 / CUDA 13.1)

Benchmark: cudf::scatter with COPYING_NVBENCH -b scatter. Baseline origin/main uses thrust::make_transform_iterator. The regressing branch uses cuda::transform_iterator over this functor:

template <typename MapType>
struct index_converter {
  __device__ MapType operator()(MapType in) const { return ((in % n_rows) + n_rows) % n_rows; }
  size_type n_rows;
};

The transformed iterator flows into thrust::scatter(..., map_it, ...), which the generic implementation expands to
thrust::copy(first, last, thrust::make_permutation_iterator(output, map_it)).

thrust::make_permutation_iterator takes the minimum traversal of its two constituents; with map_it downgraded to input, the permutation iterator becomes input, thrust::copy loses its CUB bulk fast path, and the scatter kernel falls onto a materializing fallback — costing ~170 ns/row of extra device work and extra device allocations (causing new RMM-pool OOMs at medium sizes):

num_rows	num_cols	thrust (ref)	cuda (cmp)	GPU Δ	%Diff
262 144	1	24.40 µs	70.66 µs	+46.3 µs	+189.6 %
2 097 152	1	159.34 µs	512.57 µs	+353.2 µs	+221.7 %
16 777 216	1	1.227 ms	4.062 ms	+2.835 ms	+231.1 %
262 144	8	171.12 µs	463.55 µs	+292.4 µs	+170.9 %
2 097 152	8	1.256 ms	3.414 ms	+2.158 ms	+171.9 %
16 777 216	8	10.18 ms	27.27 ms	+17.09 ms	+167.9 %

How to Reproduce

1. Check out branch in Use stream pool for gather/scatter. rapidsai/cudf#14162, but commit 2076570046683348807e34655ef9bca7fda1c9a9.
2. Build COPYING_NVBENCH.
3. COPYING_NVBENCH -b scatter --json broken.json on any NVIDIA GPU.
4. Checkout the latest commit on that branch and run .
5. COPYING_NVBENCH -b scatter --json fixed.json on any NVIDIA GPU.
6. nvbench_compare.py broken.json fixed.json

Expected behavior

There should be significant change in the performance when switching from thrust::transform_iterator vs cuda::transform_iterator.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[github-actions]

Hi @ttnghia!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[ttnghia]

Here is the affect of this bug on our code: after changing back to thrust:: iterator in rapidsai/cudf@b7c0e1f:

# scatter

## [0] Quadro RTX 6000

|  num_rows  |  num_cols  |  coalesce  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |          Diff |   %Diff |  Status  |
|------------|------------|------------|------------|-------------|------------|-------------|---------------|---------|----------|
|    4096    |     1      |     1      |   9.845 us |      84.68% |   9.008 us |      90.91% |     -0.836 us |  -8.49% |   SAME   |
|   32768    |     1      |     1      |  16.222 us |      99.85% |   9.540 us |      92.54% |     -6.682 us | -41.19% |   SAME   |
|   262144   |     1      |     1      |  70.663 us |      32.37% |  24.492 us |      20.65% |    -46.171 us | -65.34% |   FAST   |
|  2097152   |     1      |     1      | 512.565 us |       8.11% | 159.337 us |       6.84% |   -353.228 us | -68.91% |   FAST   |
|  16777216  |     1      |     1      |   4.062 ms |       2.50% |   1.227 ms |       2.21% |  -2835.560 us | -69.81% |   FAST   |
|    4096    |     8      |     1      |  50.009 us |      18.08% |  46.303 us |       8.93% |     -3.706 us |  -7.41% |   SAME   |
|   32768    |     8      |     1      |  87.459 us |      10.95% |  52.403 us |      15.24% |    -35.056 us | -40.08% |   FAST   |
|   262144   |     8      |     1      | 463.553 us |       8.59% | 149.400 us |       6.83% |   -314.152 us | -67.77% |   FAST   |
|  2097152   |     8      |     1      |   3.414 ms |       2.66% |   1.229 ms |       2.35% |  -2184.594 us | -63.99% |   FAST   |
|  16777216  |     8      |     1      |  27.271 ms |       2.87% |  10.195 ms |       4.23% | -17076.430 us | -62.62% |   FAST   |
|    4096    |    100     |     1      | 547.008 us |       3.45% | 545.607 us |       3.09% |     -1.401 us |  -0.26% |   SAME   |
|   32768    |    100     |     1      |   1.061 ms |       4.07% | 616.924 us |       1.15% |   -444.340 us | -41.87% |   FAST   |
|   262144   |    100     |     1      |   5.588 ms |       2.75% |   1.793 ms |       2.21% |  -3794.868 us | -67.91% |   FAST   |
|  2097152   |    100     |     1      |  41.344 ms |       1.72% |  16.085 ms |       4.85% | -25259.478 us | -61.10% |   FAST   |
|    4096    |    1000    |     1      |   5.606 ms |       3.38% |   5.589 ms |       4.20% |    -17.612 us |  -0.31% |   SAME   |
|   32768    |    1000    |     1      |  10.564 ms |       2.13% |   6.494 ms |       3.72% |  -4069.111 us | -38.52% |   FAST   |
|   262144   |    1000    |     1      |  55.659 ms |       2.10% |  18.660 ms |       4.12% | -36999.112 us | -66.47% |   FAST   |

[bernhardmgruber]

Thanks for reporting this issue! One difficulty is, as you pointed out, that Thrust and C++20 use different iterator models and the new CUDA iterators try to be compliant with C++20. We made several attempts to mitigate such regressions though and it seems we are not done yet. @miscco can you please take a look?