Support fancy iterators in vectorized transform and port thrust::tabulate to it by bernhardmgruber · Pull Request #6012 · NVIDIA/cccl

bernhardmgruber · 2025-09-24T18:12:14Z

thrust::sequence and thrust::tabulate before and after:

B200

# sequence

## [0] NVIDIA B200

|  T{ct}  |  Elements  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |    2^16    |   8.726 us |       4.23% |   8.950 us |       4.17% |    0.224 us |   2.57% |   SAME   |
|   I8    |    2^20    |   9.672 us |       3.85% |   9.031 us |       8.99% |   -0.641 us |  -6.63% |   FAST   |
|   I8    |    2^24    |  25.408 us |       1.22% |  11.243 us |       2.92% |  -14.165 us | -55.75% |   FAST   |
|   I8    |    2^28    | 275.036 us |       0.20% |  44.473 us |       0.95% | -230.563 us | -83.83% |   FAST   |
|   I16   |    2^16    |   8.681 us |       3.93% |   8.832 us |       4.84% |    0.151 us |   1.74% |   SAME   |
|   I16   |    2^20    |   9.707 us |       2.99% |   9.138 us |       4.58% |   -0.569 us |  -5.86% |   FAST   |
|   I16   |    2^24    |  25.654 us |       2.27% |  13.627 us |       2.73% |  -12.027 us | -46.88% |   FAST   |
|   I16   |    2^28    | 275.101 us |       0.16% |  79.732 us |       0.96% | -195.368 us | -71.02% |   FAST   |
|   I32   |    2^16    |   8.865 us |       8.48% |   9.033 us |       3.60% |    0.168 us |   1.89% |   SAME   |
|   I32   |    2^20    |   9.757 us |       4.64% |   9.463 us |       4.26% |   -0.294 us |  -3.01% |   SAME   |
|   I32   |    2^24    |  25.449 us |       1.43% |  17.660 us |       2.32% |   -7.789 us | -30.60% |   FAST   |
|   I32   |    2^28    | 275.090 us |       0.14% | 150.338 us |       0.43% | -124.752 us | -45.35% |   FAST   |
|   I64   |    2^16    |   8.701 us |       6.58% |   8.884 us |       4.88% |    0.184 us |   2.11% |   SAME   |
|   I64   |    2^20    |   9.958 us |       3.20% |  10.153 us |       3.33% |    0.195 us |   1.96% |   SAME   |
|   I64   |    2^24    |  26.378 us |       1.63% |  26.722 us |       2.78% |    0.343 us |   1.30% |   SAME   |
|   I64   |    2^28    | 291.084 us |       0.40% | 291.368 us |       0.38% |    0.284 us |   0.10% |   SAME   |

# seg_size

## [0] NVIDIA B200

|  T{ct}  |  Elements  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |        Diff |   %Diff |  Status  |
|---------|------------|------------|-------------|------------|-------------|-------------|---------|----------|
|   I8    |    2^16    |  10.114 us |       3.24% |  10.642 us |       4.50% |    0.528 us |   5.22% |   SLOW   |
|   I8    |    2^20    |  11.566 us |       2.66% |  10.818 us |       5.76% |   -0.747 us |  -6.46% |   FAST   |
|   I8    |    2^24    |  42.721 us |       1.53% |  17.983 us |       2.28% |  -24.739 us | -57.91% |   FAST   |
|   I8    |    2^28    | 541.304 us |       0.10% | 125.420 us |       0.34% | -415.884 us | -76.83% |   FAST   |
|   I16   |    2^16    |  10.094 us |       4.36% |   9.902 us |       5.06% |   -0.192 us |  -1.90% |   SAME   |
|   I16   |    2^20    |  12.110 us |       4.82% |  10.572 us |       3.11% |   -1.538 us | -12.70% |   FAST   |
|   I16   |    2^24    |  46.531 us |       2.00% |  23.409 us |       1.68% |  -23.123 us | -49.69% |   FAST   |
|   I16   |    2^28    | 584.457 us |       0.09% | 211.022 us |       0.32% | -373.435 us | -63.89% |   FAST   |
|   I32   |    2^16    |  10.137 us |       4.13% |   9.696 us |       5.77% |   -0.442 us |  -4.36% |   FAST   |
|   I32   |    2^20    |  12.144 us |       4.38% |  11.012 us |       2.85% |   -1.132 us |  -9.32% |   FAST   |
|   I32   |    2^24    |  50.881 us |       1.04% |  32.073 us |       1.54% |  -18.809 us | -36.97% |   FAST   |
|   I32   |    2^28    | 643.465 us |       0.08% | 339.741 us |       0.14% | -303.724 us | -47.20% |   FAST   |
|   I64   |    2^16    |  10.248 us |       4.25% |   9.777 us |       6.25% |   -0.470 us |  -4.59% |   FAST   |
|   I64   |    2^20    |  13.325 us |       2.16% |  12.565 us |       2.51% |   -0.761 us |  -5.71% |   FAST   |
|   I64   |    2^24    |  61.892 us |       0.68% |  49.092 us |       1.32% |  -12.800 us | -20.68% |   FAST   |
|   I64   |    2^28    | 808.201 us |       0.10% | 622.470 us |       0.20% | -185.731 us | -22.98% |   FAST   |

Fixes: #5808

copy-pr-bot · 2025-09-24T18:12:19Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

miscco · 2025-09-26T06:50:06Z

+    (int{sizeof(it_value_t<RandomAccessIteratorsIn>)}
+     * THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIteratorsIn>) ...,
+    int{size_of<it_value_t<RandomAccessIteratorOut>>});


🙀 I believe those warrant a slightly more elaborate comment

I think this needs a bigger refactoring in general

miscco · 2025-09-28T15:39:31Z

+  constexpr int element_size     = int{first_nonzero_value(
+    (sizeof(it_value_t<RandomAccessIteratorsIn>)
+     * THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIteratorsIn>) ...,
+    size_of<output_t>)};


We should really pull that out into a function

I am strongly considering to refactor the entire mess, so let's postpone any small fixes for now.

miscco · 2025-09-29T05:35:56Z

+  static constexpr bool can_memcpy_contiguous_inputs =
+    ((!THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIteratorsIn>
+      || THRUST_NS_QUALIFIER::is_trivially_relocatable_v<it_value_t<RandomAccessIteratorsIn>>)
+     && ...);


We should have a trait for that, I guess it will come up more often

github-actions · 2025-09-29T18:47:31Z

🥳 CI Workflow Results

🟩 Finished in 3h 48m: Pass: 100%/185 | Total: 8d 00h | Max: 3h 34m | Hits: 18%/189478

See results here.

…ectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Upstream: NVIDIA/cccl@ad70011 cuda tabulate.h: took upstream (drop __tabulate::functor/parallel_for; implement via cuda_cub::transform_n over ::cuda::counting_iterator with ::cuda::proclaim_copyable_arguments(tabulate_op)); fancy-iterator support itself is in cub DeviceTransform. HIP: no change: performance-only (system/hip tabulate keeps its parallel_for path). benchmark/: bench/tabulate/basic.cu adapted to primbench: added sequence_benchmark (subalgo "sequence"), seg_size_t static_cast, output vectors use thrust::no_init, types widened to int8/16/32/64 (upstream integral_types); kept subalgo "basic" for the seg_size benchmark (upstream renamed it seg_size). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Upstream: NVIDIA/cccl@ad70011 cuda tabulate.h: took upstream (drop __tabulate::functor/parallel_for; implement via cuda_cub::transform_n over ::cuda::counting_iterator with ::cuda::proclaim_copyable_arguments(tabulate_op)); fancy-iterator support itself is in cub DeviceTransform. HIP: no change: performance-only (system/hip tabulate keeps its parallel_for path). benchmark/: bench/tabulate/basic.cu adapted to primbench: added sequence_benchmark (subalgo "sequence"), seg_size_t static_cast, output vectors use thrust::no_init, types widened to int8/16/32/64 (upstream integral_types); kept subalgo "basic" for the seg_size benchmark (upstream renamed it seg_size). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…ectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Upstream: NVIDIA/cccl@ad70011 cuda tabulate.h: took upstream (drop __tabulate::functor/parallel_for; implement via cuda_cub::transform_n over ::cuda::counting_iterator with ::cuda::proclaim_copyable_arguments(tabulate_op)); fancy-iterator support itself is in cub DeviceTransform. HIP: no change: performance-only (system/hip tabulate keeps its parallel_for path). benchmark/: bench/tabulate/basic.cu adapted to primbench: added sequence_benchmark (subalgo "sequence"), seg_size_t static_cast, output vectors use thrust::no_init, types widened to int8/16/32/64 (upstream integral_types); kept subalgo "basic" for the seg_size benchmark (upstream renamed it seg_size). test/: no counterpart. testing/: no counterpart. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…sform over counting_iterator Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Related local port commit: 291f437 The local port took upstream's cuda tabulate.h but left system/hip on its own parallel_for functor as "no change: performance-only". Bring system/hip in line with upstream's structure: - Drop __tabulate::functor; tabulate is a transform of _THRUST_LIBCXX::counting_iterator<iter_difference_t<Iterator>> into the output range, with granular libhipcxx includes as upstream. - hip has no transform_n overload (NVIDIA/cccl#5238 was cuda-only), so hip_rocprim::transform is called with [begin, begin + count). - tabulate_op is not wrapped in _THRUST_LIBCXX::proclaim_copyable_arguments: the proclamation only feeds cub::DeviceTransform, which hip's transform does not use, and the wrapper inherits from the functor, so it would reject function pointers and final classes for no benefit. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>

…sform_n with copyable arguments Related upstream commit: ad70011e990 Support fancy iterators in vectorized transform and port thrust::tabulate to it (NVIDIA/cccl#6012) Upstream: NVIDIA/cccl@ad70011 Related local port commit: 39a06d7 Related local parity commit: 4698506 4698506 turned hip tabulate into a transform over a counting_iterator, but called transform(begin, begin + count, ...) with the bare tabulate_op. Finish matching upstream in thrust/system/hip/detail/tabulate.h: call hip_rocprim::transform_n (added by the 76fc4572271 parity commit) on _THRUST_LIBCXX::counting_iterator<size_type>{} and count, and wrap tabulate_op in _THRUST_LIBCXX::proclaim_copyable_arguments. Includes follow upstream: __functional/address_stability.h and the _THRUST_STD iterator header replace __iterator/distance.h and __iterator/incrementable_traits.h. The counting_iterator include stays, since upstream only gets it through transform.h. Not compiled, tests not run Co-Authored-By: Claude <noreply@anthropic.com>

github-project-automation Bot added this to CCCL Sep 24, 2025

github-project-automation Bot moved this to Todo in CCCL Sep 24, 2025

cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Sep 24, 2025

bernhardmgruber force-pushed the fancy_transform branch from 2208f23 to d69fd95 Compare September 25, 2025 18:05

miscco reviewed Sep 26, 2025

View reviewed changes

bernhardmgruber force-pushed the fancy_transform branch 2 times, most recently from 3a9e091 to 512baf5 Compare September 26, 2025 10:50

bernhardmgruber marked this pull request as ready for review September 26, 2025 11:05

bernhardmgruber requested review from a team as code owners September 26, 2025 11:05

bernhardmgruber requested review from alliepiper and shwina September 26, 2025 11:05

cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Sep 26, 2025

NaderAlAwar reviewed Sep 26, 2025

View reviewed changes

Comment thread thrust/benchmarks/bench/tabulate/basic.cu Resolved

Comment thread cub/cub/device/dispatch/kernels/transform.cuh Resolved

This comment has been minimized.

Sign in to view

bernhardmgruber force-pushed the fancy_transform branch from 81f3ede to 8d22b69 Compare September 27, 2025 03:43

This comment has been minimized.

Sign in to view

bernhardmgruber changed the title ~~Support mixing fancy iterators with normal ones in vectorized transform~~ Support fancy iterators in vectorized transform and port thrust::tabulate to it Sep 28, 2025

miscco approved these changes Sep 29, 2025

View reviewed changes

bernhardmgruber added 9 commits September 29, 2025 16:54

Add benchmark for thrust::sequence

6a14dbf

Extend benchmark to all integer types

9644377

Support mixing fancy iterators with normal ones in vectorized transform

5c48e92

Port thrust::tabulate to thrust::transform

e1d6ed9

Fix

ce02b53

Improve clarity

6543c30

fix

2fbf093

Drop obsolete comment

c13c974

Reviewer feedback

c2cb534

bernhardmgruber force-pushed the fancy_transform branch from 8d22b69 to c2cb534 Compare September 29, 2025 14:54

bernhardmgruber enabled auto-merge (squash) September 29, 2025 19:14

NaderAlAwar approved these changes Sep 30, 2025

View reviewed changes

bernhardmgruber merged commit ad70011 into NVIDIA:main Sep 30, 2025
1 check passed

github-project-automation Bot moved this from In Review to Done in CCCL Sep 30, 2025

bernhardmgruber deleted the fancy_transform branch September 30, 2025 13:00

This was referenced May 19, 2026

thrust::tabulate is much slower than equivalent transform/for_each for wide row functor #9070

Open

Avoid copyable-argument transform path in tabulate #9071

Closed

gevtushenko mentioned this pull request Sep 7, 2026

[EPIC]: Regression-prevention skills: perf, functional, API, ABI #11244

Closed

1 task done

Navigation Menu

Repository navigation

Support fancy iterators in vectorized transform and port thrust::tabulate to it - #6012

Support fancy iterators in vectorized transform and port thrust::tabulate to it#6012
bernhardmgruber merged 9 commits into
NVIDIA:mainfrom
bernhardmgruber:fancy_transform

bernhardmgruber commented Sep 24, 2025 •

edited

Loading

Uh oh!

copy-pr-bot Bot commented Sep 24, 2025

Uh oh!

miscco Sep 26, 2025

Uh oh!

bernhardmgruber Sep 26, 2025

Uh oh!

bernhardmgruber Sep 26, 2025

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

This comment has been minimized.

miscco Sep 28, 2025

Uh oh!

bernhardmgruber Sep 29, 2025

Uh oh!

Uh oh!

miscco Sep 29, 2025

Uh oh!

Uh oh!

Uh oh!

github-actions Bot commented Sep 29, 2025

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

bernhardmgruber commented Sep 24, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

copy-pr-bot Bot commented Sep 24, 2025

Uh oh!

miscco Sep 26, 2025

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Sep 26, 2025

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Sep 26, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

This comment has been minimized.

This comment has been minimized.

miscco Sep 28, 2025

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Sep 29, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

miscco Sep 29, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

github-actions Bot commented Sep 29, 2025

🥳 CI Workflow Results

🟩 Finished in 3h 48m: Pass: 100%/185 | Total: 8d 00h | Max: 3h 34m | Hits: 18%/189478

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

bernhardmgruber commented Sep 24, 2025 •

edited

Loading