Repository navigation
Conversation
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:
No actionable comments were generated in the recent review.
 |
| Layer / File(s) | Summary |
|---|---|
Add striped vectorized loading cub/cub/block/block_load.cuh |
Adds the algorithm enum, loader helpers, unguarded dispatch, and guarded-load routing. |
Use vectorized loads in CUB algorithms cub/cub/agent/agent_find.cuh, cub/cub/agent/agent_reduce.cuh, cub/cub/device/dispatch/kernels/kernel_transform.cuh |
Find, reduce, and transform use the loading helpers in vectorized paths. Reduce derives its vector type from the effective load width. |
Test striped vectorized loads cub/test/catch2_test_block_load.cu, cub/test/catch2_test_enum_formatting.cu |
Tests cover striped-vectorized indexing, counting iterators, alignment, conversion cases, and enum formatting. |
HyperLogLog Vectorized Access
| Layer / File(s) | Summary |
|---|---|
Update HyperLogLog vectorized access cudax/include/cuda/experimental/__cuco/detail/hyperloglog/hyperloglog_impl.cuh, cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh |
Alignment calculation uses the maximum vector-access limit. The kernel uses byte-vectorized storage and element aliases. |
Possibly related issues
- NVIDIA/cccl#12170 — Requests the striped vectorized load algorithm added here.
Priority:
Change: Feature
Merge Risk: 
bac3d
-
Aligned custom types can take a load path that bypasses their normal per-element copy behavior. Restrict that path before merging; convertible input types are supported.
Comment @coderabbitai help to get the list of available commands.
This comment has been minimized.
This comment has been minimized.
fbusato
reviewed
Oct 16, 2026
fbusato
left a comment
fbusato
left a comment
Contributor
There was a problem hiding this comment.
this makes me think that we are keeping reinventing the wheel, as we could use cuda::std::simd directly
|
|
||
| #ifndef _CCCL_DOXYGEN_INVOKED // Do not document | ||
|
|
||
| //! @brief Internal implementation of LoadDirectStripedVectorized for input aligned to the vector type |
Contributor
There was a problem hiding this comment.
if it is internal, this should be placed in the detail namespace
| // Vector i of a thread is the blocked vector of thread linear_tid + i * ThreadsPerBlock | ||
| // The thread offset is computed in unsigned arithmetic like offsets from threadIdx.x, because a signed multiplication | ||
| // generates additional address arithmetic in the calling kernels | ||
| const T* thread_src_ptr = block_src_ptr + static_cast<unsigned int>(linear_tid) * VecSize; |
Contributor
There was a problem hiding this comment.
Suggested change
| const T* thread_src_ptr = block_src_ptr + static_cast<unsigned int>(linear_tid) * VecSize; | |
| const T* thread_src_ptr = block_src_ptr + static_cast<unsigned>(linear_tid) * VecSize; |
we could also use _CCCL_ASSUME(linear_tid >= 0)
| InternalLoadDirectStripedVectorized(int linear_tid, const T* block_src_ptr, T (&dst_items)[ItemsPerThread]) | ||
| { | ||
| static_assert(ItemsPerThread % VecSize == 0, "VecSize must evenly divide ItemsPerThread"); | ||
|
|
Contributor
There was a problem hiding this comment.
we also need to check the actual alignment
Comment on lines
+352
to
+357
| _CCCL_PRAGMA_UNROLL_FULL() | ||
| for (int i = 0; i < ItemsPerThread / VecSize; i++) | ||
| { | ||
| InternalLoadDirectBlockedVectorized<Modifier>( | ||
| i * ThreadsPerBlock, thread_src_ptr, reinterpret_cast<T(&)[VecSize]>(dst_items[i * VecSize])); | ||
| } |
Contributor
There was a problem hiding this comment.
(I didn't compile it)
Suggested change
| _CCCL_PRAGMA_UNROLL_FULL() | |
| for (int i = 0; i < ItemsPerThread / VecSize; i++) | |
| { | |
| InternalLoadDirectBlockedVectorized<Modifier>( | |
| i * ThreadsPerBlock, thread_src_ptr, reinterpret_cast<T(&)[VecSize]>(dst_items[i * VecSize])); | |
| } | |
| constexpr auto num_vecs = ItemsPerThread / VecSize; | |
| auto& dst_vecs = reinterpret_cast<T(*)[VecSize][num_vecs]>(dst_items); | |
| _CCCL_PRAGMA_UNROLL_FULL() | |
| for (int i = 0; i < ItemsPerThread / VecSize; i++) | |
| { | |
| InternalLoadDirectBlockedVectorized<Modifier>(i * ThreadsPerBlock, thread_src_ptr, dst_vecs[i]); | |
| } |
| // Vectors of the width BLOCK_LOAD_VECTORIZE loads, or single items if that does not hold a whole number of T | ||
| using vector_word_t = typename ::cuda::__as_bytes_vectorized_t<T, ItemsPerThread>::value_type; | ||
| constexpr int vec_size = | ||
| sizeof(vector_word_t) % sizeof(T) == 0 ? static_cast<int>(sizeof(vector_word_t) / sizeof(T)) : 1; |
Contributor
There was a problem hiding this comment.
Suggested change
| sizeof(vector_word_t) % sizeof(T) == 0 ? static_cast<int>(sizeof(vector_word_t) / sizeof(T)) : 1; | |
| sizeof(vector_word_t) % sizeof(T) == 0 ? int{sizeof(vector_word_t) / sizeof(T)} : 1; |
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
block-load-striped-vectorize
branch
from
October 19, 2026 09:38
478dea5 to
144aee2
Compare
![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: 2
ℹ️ Review info
⚙️ Run configuration
- Configuration used: Repository: NVIDIA/cccl/.coderabbit.yaml
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
69290ac0-8724-4216-b330-2dd31a170789
📒 Files selected for processing: 4
cub/cub/agent/agent_find.cuhcub/cub/agent/agent_reduce.cuhcub/cub/block/block_load.cuhcub/cub/device/dispatch/kernels/kernel_transform.cuh
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 10 remain after this review.
Comment on lines
+422
to
+424
| if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator> && ::cuda::__is_vectorizable<T>()) | ||
| { | ||
| const T* block_src_ptr = ::cuda::std::to_address(block_src_it); |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
important: Keep convertible input types on the item-by-item path. If BlockLoad<float, ..., BLOCK_LOAD_STRIPED_VECTORIZED>::Load receives an int*, the contiguous branch tries to assign that pointer to const float* and fails to compile. BlockLoad documents that input values need only be convertible to T. Require the iterator value type to be T before selecting vector loads; otherwise use the existing item-by-item path. As per path instructions: “Focus on algorithm correctness, temporary-storage protocol, dispatch/policy selection, stream behavior, CUDA error handling, synchronization, memory access safety, performance regressions, and test coverage.”
Source: Path instructions
Comment on lines
+422
to
+428
| if constexpr (THRUST_NS_QUALIFIER::is_contiguous_iterator_v<RandomAccessIterator> && ::cuda::__is_vectorizable<T>()) | ||
| { | ||
| const T* block_src_ptr = ::cuda::std::to_address(block_src_it); | ||
| // Add the alignment check to ensure the vectorized loading can proceed. | ||
| if (::cuda::std::is_sufficiently_aligned<alignof(::cuda::__as_bytes_vectorized_t<T, VecSize>)>(block_src_ptr)) | ||
| { | ||
| InternalLoadDirectStripedVectorized<ThreadsPerBlock, VecSize, Modifier>(linear_tid, block_src_ptr, dst_items); |
Contributor
There was a problem hiding this comment.
🎯 Functional Correctness | 🟠 Major | ⚡ Quick win
important: Do not byte-load types with nontrivial copy behavior. __is_vectorizable<T>() checks alignment only. For a suitably aligned T with a nontrivial copy operation, this branch loads its representation into vector storage and aliases that storage as T rather than loading a T object. The new algorithm can therefore bypass the type’s copy behavior. Restrict this branch to types that support representation-based copying, and use the item-by-item path otherwise. (raw.githubusercontent.com) As per path instructions: “Focus on algorithm correctness, temporary-storage protocol, dispatch/policy selection, stream behavior, CUDA error handling, synchronization, memory access safety, performance regressions, and test coverage.”
Source: Path instructions
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
block-load-striped-vectorize
branch
from
October 23, 2026 12:36
144aee2 to
25c09c7
Compare
miscco
removed this pull request from stack #12044
October 23, 2026 12:36
miscco
changed the base branch from
as_vectorized_bytes
to
vectorized_device_for
October 23, 2026 12:39
miscco
added this pull request to stack #12219
October 23, 2026 12:39
miscco
force-pushed
the
block-load-striped-vectorize
branch
from
October 23, 2026 12:43
25c09c7 to
26bf61c
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
block-load-striped-vectorize
branch
2 times, most recently
from
October 24, 2026 17:35
10d959f to
cbe8014
Compare
Contributor
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to GitHub limitations.
kernels.cuh:88-95
cudax/include/cuda/experimental/__cuco/detail/hyperloglog/kernels.cuh:88-95
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick winimportant: Preserve typed copies for non-trivially copyable
_Tp.
add_asyncaccepts convertible value types, and a contiguousT*reaches__add_shmem_vectorized. The new byte storage copies words, then__alias<_Tp>accesses them without constructing or invoking_Tp's copy operations. This bypasses the typed element copies used by the previouscuda::std::array<_Tp, _VectorSize>path and can produce invalid objects or incorrect hash updates.Suggested fix
-
-
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: vectorized&
#95;device_for - File: .coderabbit.yaml
- Commit: 53436ab
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
f1e40a8d-63cd-47a9-863c-0057c9a2db2a
-
cub/cub/block/block_load.cuhcub/test/catch2_test_block_load.cu
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 9 remain after this review.
-
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
miscco
force-pushed
the
block-load-striped-vectorize
branch
from
October 26, 2026 09:25
cbe8014 to
99bf067
Compare
Load vectors of `cuda::__as_bytes_vectorized_t` instead of `cuda::std::array` and access the elements through `__alias`. The widest vector access is taken from `cuda::__max_vector_access`, which is 16 bytes before CTK 13.0, instead of being hard coded to 32 bytes.
This adds machinery for striped vectorized loads to `CUB`. There are multiple places were we need to load vectors in a striped arrangement, for example in `DeviceReduce`, `DeviceFind` and the vectorized `DeviceTransform` kernel. Having this open coded complicates maintenance and makes it harder to reason about the code. In the future we want to add iterator unwrapping to extend the vectorized loads to non-contiguous inputs. `LoadDirectStripedVectorized` loads a striped arrangement of vectors of `VecSize` consecutive items: thread `t` loads the vectors `t`, `t + ThreadsPerBlock`, and so on. Like `LoadDirectBlockedVectorized`, it loads `__as_bytes_vectorized_t` vectors of the value type of the input if the iterator is contiguous, that type is vectorizable and the input is aligned, and otherwise loads the same arrangement item by item. The items are converted to `T` when they are copied out of the vectors. `InternalLoadDirectStripedVectorized` only performs the vectorized loads and asserts that the input is aligned. `BLOCK_LOAD_STRIPED_VECTORIZED` uses vectors of the width that `BLOCK_LOAD_VECTORIZE` loads. Guarded loads use the arrangement of `BLOCK_LOAD_STRIPED`. `DeviceReduce`, `DeviceFind` and the vectorized `DeviceTransform` kernel load their full tiles with `InternalLoadDirectStripedVectorized`, and `DeviceTransform` loads non-contiguous inputs with `LoadDirectStripedVectorized`. The `VectorT` of `AgentReduce` is now sized by `vec_size` instead of `VECTOR_LOAD_LENGTH`, so policies with fewer items per thread than `VECTOR_LOAD_LENGTH` no longer load vectors wider than their items.
miscco
force-pushed
the
block-load-striped-vectorize
branch
from
October 26, 2026 09:26
99bf067 to
bac3d39
Compare
Contributor
🔬 Benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Contributor
😬 CI Workflow Results🟥 Finished in 3h 29m: Pass: 99%/254 | Total: 13d 01h | Max: 3h 26m | Hits: 22%/1579109See results here. AI failure analysis1. CUB WarpExchangeShfl lane_id field shadowing under Clang 21 · 1 jobExplanation: The new `kernel_transform.cuh` dependency on `cub/block/block_load.cuh` exposes `warp_exchange_shfl.cuh` to libcudacxx mdspan compilations. Clang 21 treats the `transpose` parameter shadowing the class field as an error because `-Wshadow` and `-Werror` are enabled, producing the same diagnostic across 122 lit precompile failures. Evidence: Copy this prompt into a coding agentJobs: |
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.
This adds machinery for striped vectorized loads to CUB. There are multiple places were we need to load vectors in a striped arrangement, for example in DeviceReduce, DeviceFind and the vectorized DeviceTransform kernel.
Having this open coded complicates maintenance and makes it harder to reason about the code.
In the future we want to add iterator unwrapping to extend the vectorized loads to non-contiguous inputs.
LoadDirectStripedVectorized loads a striped arrangement of vectors of VecSize consecutive items: thread t loads
the vectors t, t + ThreadsPerBlock, and so on. Like LoadDirectBlockedVectorized, it loads __as_bytes_vectorized_t
vectors if the iterator is contiguous, the type is vectorizable and the input is aligned, and otherwise loads the same
arrangement item by item. InternalLoadDirectStripedVectorized only performs the vectorized loads and asserts that the
input is aligned.
BLOCK_LOAD_STRIPED_VECTORIZED uses vectors of the width that BLOCK_LOAD_VECTORIZE loads. Guarded loads use the
arrangement of BLOCK_LOAD_STRIPED.
DeviceReduce, DeviceFind and the vectorized DeviceTransform kernel load their full tiles with
InternalLoadDirectStripedVectorized, and DeviceTransform loads non-contiguous inputs with
LoadDirectStripedVectorized. The VectorT of AgentReduce is now sized by vec_size instead of VECTOR_LOAD_LENGTH,
so policies with fewer items per thread than VECTOR_LOAD_LENGTH no longer load vectors wider than their items.
Stack created with GitHub Stacks CLI • Give Feedback 💬