Repository navigation
Conversation
miscco
requested review from
griwes,
shwina and
sleeepyjack
and removed request for
a team
October 16, 2026 18:48
miscco
requested review from
elstehle and
gonidelis
and removed request for
a team
October 16, 2026 18:48
Contributor
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
📝 Summary
Priority: Change: Feature Merge Risk:
 |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Actionable comments posted: 4
ℹ️ Review info
⚙️ Run configuration
- Configuration used: Repository: NVIDIA/cccl/.coderabbit.yaml
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
b8707118-c732-4dc0-9abc-6e69372472fd
📒 Files selected for processing: 35
cub/benchmarks/bench/load/block_load.cucub/benchmarks/bench/reduce/contiguous_adaptor.cucub/benchmarks/bench/topk/transform_iterator.cucub/benchmarks/bench/transform/contiguous_adaptor.cucub/cub/agent/agent_find.cuhcub/cub/agent/agent_reduce.cuhcub/cub/block/block_load.cuhcub/cub/device/device_for.cuhcub/cub/device/dispatch/dispatch_transform.cuhcub/cub/device/dispatch/kernels/kernel_transform.cuhcub/test/catch2_test_block_load.cucub/test/catch2_test_device_reduce_iterators.cucub/test/catch2_test_device_transform_vectorized.cucub/test/catch2_test_enum_formatting.cucudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuhlibcudacxx/include/cuda/__iterator/contiguous_iterator_adaptor.hlibcudacxx/include/cuda/__iterator/transform_input_output_iterator.hlibcudacxx/include/cuda/__iterator/transform_iterator.hlibcudacxx/include/cuda/__iterator/transform_output_iterator.hlibcudacxx/include/cuda/__memory/as_bytes_vectorized.hlibcudacxx/include/cuda/iteratorlibcudacxx/include/cuda/std/__iterator/iterator_traits.hlibcudacxx/test/libcudacxx/cuda/iterators/contiguous_iterator_adaptor.pass.cpplibcudacxx/test/libcudacxx/libcxx/iterators/has_contiguous_traversal.compile.pass.cpplibcudacxx/test/libcudacxx/libcxx/memory/as_bytes_vectorized.compile.pass.cppthrust/testing/contiguous_iterator_adaptor.cuthrust/testing/is_contiguous_iterator.cuthrust/testing/trivial_sequence.cuthrust/testing/type_traits.cuthrust/thrust/detail/pointer.hthrust/thrust/device_ptr.hthrust/thrust/iterator/detail/normal_iterator.hthrust/thrust/iterator/transform_input_output_iterator.hthrust/thrust/iterator/transform_iterator.hthrust/thrust/iterator/transform_output_iterator.h
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.
Comment on lines
+71
to
+73
| if (threadIdx.x == 0) | ||
| { | ||
| block_sums[blockIdx.x] = sum; |
Contributor
There was a problem hiding this comment.
🚀 Performance & Scalability | 🟠 Major | 🏗️ Heavy lift
important: Make every thread's loaded data affect the result.
Only thread 0 writes its sum. For both direct and vectorized loads, the other threads' loaded values have no observable use, so the compiler can remove their loads. The benchmark then reports num_items reads for work that may read only one thread's items per tile. Reduce the per-thread sums across the block before writing block_sums. CUB's direct and vectorized BlockLoad paths use thread-local loads. (raw.githubusercontent.com)
As per path instructions, benchmark changes must “measure meaningful workloads” and “preserve useful comparison data.”
Source: Path instructions
| template <class Env> | ||
| void max_keys(nvbench::state& state, int64_t elements, int selected_elements, Env env) | ||
| { | ||
| thrust::device_vector<char> input(static_cast<size_t>(elements), thrust::no_init); |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Initialize the TopK input keys.
thrust::no_init leaves input uninitialized, and this benchmark never writes the keys before MaxKeys reads them. The TopK workload therefore depends on prior allocation contents instead of a reproducible key distribution. Fill input with a controlled distribution before measuring either load policy. (nvidia.github.io)
As per path instructions, benchmark changes must “measure meaningful workloads” and “preserve useful comparison data.”
Source: Path instructions
| state.add_global_memory_writes<uint32_t>(selected_elements, "OutputKeys"); | ||
|
|
||
| size_t temp_size{}; | ||
| cub::DeviceTopK::MaxKeys(nullptr, temp_size, in, out, elements, selected_elements, env); |
Contributor
There was a problem hiding this comment.
🩺 Stability & Availability | 🟡 Minor | ⚡ Quick win
suggestion: Check both MaxKeys results.
MaxKeys returns cudaError_t, but the storage query and measured call discard it. If either call fails, the benchmark can continue without reporting the CUB failure. Apply the existing runtime-API error wrapper to both calls. (raw.githubusercontent.com)
As per path instructions, cub changes require attention to “CUDA error handling.”
Also applies to: 71-71
Source: Path instructions
Comment on lines
+288
to
+295
| using value_t = it_value_t<decltype(in)>; | ||
| if constexpr (::cuda::__is_contiguous_iterator_or_adaptor_v<decltype(in)>) | ||
| { // CanVectorize only selects this path if contiguous inputs are aligned to their vectors | ||
| const auto* base_ptr = ::cuda::__unwrap_iterator_adaptor(in); | ||
| _CCCL_ASSUME(reinterpret_cast<::cuda::std::uintptr_t>(base_ptr) % (sizeof(value_t) * VecSize) == 0); | ||
| } | ||
| auto& items = *reinterpret_cast<value_t(*)[ItemsPerThread]>(input.data()); | ||
| LoadDirectStripedVectorized<ThreadsPerBlock, VecSize>(static_cast<int>(threadIdx.x), in, items); |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: DeviceTransform computes vector alignment from the adaptor's value type. The kernel loads vectors of the unwrapped storage type. If the storage type is larger than the value type, _CCCL_ASSUME asserts an alignment that the host never checked. The compiler can then emit misaligned vector loads. If the storage type is smaller, vectorization is skipped more often than needed.
cub/cub/device/dispatch/kernels/kernel_transform.cuh#L288-L295: compute the assumed alignment assizeof(::cuda::__unwrapped_iter_value_t<decltype(in)>) * VecSize.cub/cub/device/dispatch/dispatch_transform.cuh#L151-L163: inCanVectorize, passsizeof(::cuda::__unwrapped_iter_value_t<In>) * vec_sizefor the input alignment check.
📍 Affects 2 files
cub/cub/device/dispatch/kernels/kernel_transform.cuh#L288-L295(this comment)cub/cub/device/dispatch/dispatch_transform.cuh#L151-L163
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 16, 2026 19:08
02bb6cb to
f0e4283
Compare
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 19, 2026 09:38
f0e4283 to
9f6ed02
Compare
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 23, 2026 12:36
9f6ed02 to
93e7607
Compare
This was referenced Oct 24, 2026
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 26, 2026 09:25
699f213 to
f577e76
Compare
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 26, 2026 09:26
f577e76 to
65f4995
Compare
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 27, 2026 20:12
65f4995 to
b168327
Compare
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 27, 2026 21:23
b168327 to
49ef269
Compare
This comment has been minimized.
This comment has been minimized.
Contributor
😬 CI Workflow Results🟥 Finished in 4h 15m: Pass: 92%/394 | Total: 16d 23h | Max: 4h 07m | Hits: 42%/3503941See results here. AI failure analysis1. CUB transform alignment helper leaves it unused · 24 jobsExplanation: The changed is_pointer_aligned template only references it in the contiguous-iterator constexpr branch. Instantiations selecting the fallback branch therefore trigger unused-parameter warnings promoted to errors across GCC and MSVC builds. Evidence: Copy this prompt into a coding agentJobs:
2. CUB WarpExchangeShfl transpose parameter shadows lane_id field · 3 jobsExplanation: The transpose function parameter named lane_id shadows the class field with the same name. ClangCUDA enables -Wshadow and -Werror, so every libcudacxx test including this header fails during precompilation. Evidence: Copy this prompt into a coding agentJobs: 3. Thrust contiguous-adaptor hooks violate clang-tidy naming rules · 1 jobExplanation: New libcudacxx-style identifiers were added inside Thrust files, whose clang-tidy configuration requires CamelCase template parameters and lower_case functions. The earliest failures are in pointer.h, with equivalent findings on the other newly added adaptor hooks and trait specializations. Evidence: Copy this prompt into a coding agentJobs: 4. Thrust internal contiguous traits break Doxygen symbol resolution · 1 jobExplanation: Doxygen attempts to document the newly added internal variable-template specializations but cannot match their macro-expanded declarations to defined symbols. These implementation traits are not public API and should be excluded using the repository's established internal-documentation pattern. Evidence: Copy this prompt into a coding agentJobs: 5. pytest-benchmark 5.3.0 accesses removed pytest Function._nodeid · 3 jobsExplanation: All three Python 3.14 example lanes install pytest 9.2.0 with pytest-benchmark 5.3.0. The benchmark fixture accesses pytest's removed private Function._nodeid attribute during setup, so all 125 host benchmark cases error before running. Evidence: Copy this prompt into a coding agentJobs: 6. SASS changes lack required approval · 1 jobExplanation: The approval artifact for this exact commit was absent, so the mandatory SASS-change gate failed. This is an independent approval requirement rather than a derivative workflow failure. Evidence: Copy this prompt into a coding agentJobs: |
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
2 times, most recently
from
October 28, 2026 13:41
4a15c4d to
c6e60f7
Compare
Contributor
🔬 Benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Fancy iterators over a pointer are not contiguous, so algorithms miss the fast copy. Unwrap peels transform, transform_output, and transform_input_output iterators to that pointer, and rebase retargets them at a copy of the storage without dropping their functors. A contiguous iterator unwraps to its address, even if it also has a base. Transform and transform_input_output iterators invoke their functor on the copy, so they can only be rebased if that functor proclaims copyable arguments.
Device pointers and vector iterators are contiguous storage. Teach them to unwrap and rebase so a transform, transform_output, or transform_input_output iterator over them can be retargeted at a copy of that storage. Like in libcu++, transform and transform_input_output iterators can only be rebased if the functor that reads the storage proclaims copyable arguments.
The vectorized block loads take an iterator and load the storage behind it with vectors into a thread-local copy. A contiguous iterator reads its items from that copy, converting them if the item type differs, like a base of the stored type. An iterator adaptor over contiguous storage, like `cuda::transform_iterator` over a pointer, reads the items through itself rebased onto the copy, so the functor runs on the loaded storage. Adaptors whose functor does not proclaim copyable arguments cannot be rebased and are loaded item by item. DeviceReduce vectorizes such adaptors, and the vectorized DeviceTransform kernel loads all of its inputs through LoadDirectStripedVectorized. Dispatch already checks the alignment of contiguous inputs, so the kernel assumes it and the compiler drops their alignment check and fallback.
miscco
force-pushed
the
unwrapped_copy_bulk_async
branch
from
October 28, 2026 14:44
c6e60f7 to
8ee66a5
Compare
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
We have a lot of places where we vectorize loading data if the input is contiguous and laso sufficiently aligned
However, we can also extend that to fancy iterators that are wrappers around such iterators and then rebase after the vectorized load.
Stack created with GitHub Stacks CLI • Give Feedback 💬