[bdice]

Summary

• Removes cuda::proclaim_copyable_arguments from the CUDA thrust::tabulate backend.
• Keeps tabulate implemented through cuda_cub::transform_n, but lets transform dispatch choose the stable-address path for user functors.
• Fixes a large slowdown found in wide row-wise tabulate workloads.

Details

In #6012, CUDA tabulate was ported from a custom parallel_for implementation to transform_n and the functor was wrapped with cuda::proclaim_copyable_arguments.

A standalone CCCL + NVBench reproducer in #9070 showed that the wrapper is the source of the slowdown. Applying proclaim_copyable_arguments manually to an otherwise fast transform_n reproduces the slow tabulate behavior.

Representative results on NVIDIA GB10, CUDA 13.2, CCCL main fcd4ff4ea1:

rows	cols	tabulate	transform_n_ptrdiff	transform_n_proclaim
1,048,576	32	4.197 ms	1.185 ms	4.277 ms
1,048,576	128	18.337 ms	4.679 ms	18.361 ms

With this patch:

rows	cols	patched tabulate	transform_n_ptrdiff	for_each	cub_for	direct_kernel
1,048,576	32	1.199 ms	1.182 ms	1.172 ms	1.191 ms	1.198 ms
1,048,576	128	4.733 ms	4.657 ms	4.558 ms	4.641 ms	4.573 ms
1,048,576	256	9.249 ms	9.312 ms	9.192 ms	9.148 ms	9.198 ms

Fixes #9070.

[coderabbitai]

📝 Walkthrough

Summary by CodeRabbit

• Chores
  • Refined internal GPU computation implementation by optimizing operation invocation patterns, with no impact to public interfaces or user-facing functionality.

Walkthrough

The PR removes a proclaim_copyable_arguments wrapper from the functor passed to transform_n in the CUDA implementation of thrust::tabulate. This aligns the internal implementation with how thrust::transform passes functors, addressing a significant performance gap between the two algorithms for wide-row functors.

Changes

thrust::tabulate CUDA implementation optimization

Layer / File(s)	Summary
tabulate transform_n functor invocation thrust/thrust/system/cuda/detail/tabulate.h	Removed ::cuda::proclaim_copyable_arguments() wrapper from the tabulate_op functor passed to transform_n, passing the functor directly to match the pattern used by thrust::transform.

Assessment against linked issues

Objective	Addressed	Explanation
Improve thrust::tabulate performance for wide-row functors to match thrust::transform, thrust::for_each_n, and cub::DeviceFor::ForEach [#9070]	✅

---

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[coderabbitai]

🧹 Nitpick comments (1)

thrust/thrust/system/cuda/detail/tabulate.h (1)

35-35: ⚡ Quick win

suggestion: The unused include at line 22 (<cuda/__functional/address_stability.h>) can be removed—transform.h already includes it. Tests cover standard functor cases (identity, placeholder expressions), but don't explicitly exercise capturing lambdas or stateful functors; if such cases are thrust::tabulate use cases, consider adding tests to verify transform_n's dispatch handles them correctly without the removed wrapper.

---

ℹ️ Review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 23b98d62-9862-4aaf-90d8-1cfb40a29aca

📥 Commits

Reviewing files that changed from the base of the PR and between fcd4ff4 and 87cfeaa.

📒 Files selected for processing (1)

• thrust/thrust/system/cuda/detail/tabulate.h

[bernhardmgruber]

I am sorry proclaim_copyable_arguments is causing a regression on your side. However, the proposed change causes a >5x regression for other workloads, see #6012 which shows benchmarks on the initial improvements.

[bdice]

Thanks for looking into it with me. I'll report some additional benchmarks on #9070, closing this for now.

[github-actions]

🥳 CI Workflow Results

🟩 Finished in 2h 21m: Pass: 100%/117 | Total: 3d 15h | Max: 2h 03m | Hits: 42%/549083

See results here.