Repository navigation
Conversation
edc0774 to
1fec506
Compare
|
Important Review skippedReview was skipped as selected files did not have any reviewable changes.
 |
| Layer / File(s) | Summary |
|---|---|
Typed vectorized loader and fallback cub/cub/block/block_load.cuh, cub/test/catch2_test_block_load.cu |
LoadDirectBlockedVectorized accepts separate input and destination types, a cache modifier, and an optional vector size. It vectorizes eligible aligned input and otherwise uses direct blocked loading. Tests cover configured types, input offsets, and vector sizes. |
Warp vectorized-load integration cub/cub/warp/warp_load.cuh |
Warp pointer and cache-modified iterator overloads call LoadDirectBlockedVectorized. The iterator overload forwards its modifier. |
TopK vector-size policy and loading cub/cub/device/dispatch/tuning/tuning_topk.cuh, cub/cub/device/dispatch/dispatch_topk.cuh, cub/cub/agent/agent_topk.cuh |
TopK policies store and report vector sizes, and dispatch passes the applicable size to each agent. Full-tile loads use direct vectorized loading when the policy selects vectorization and the input iterator is contiguous; otherwise they use BlockLoad. |
Possibly related PRs
- NVIDIA/cccl#12049: Both PRs change blocked vectorized-load eligibility and handling of differing input types.
- NVIDIA/cccl#12163: Both PRs modify type and volatility checks for
BlockLoadvectorized dispatch. - NVIDIA/cccl#12042: This PR uses the byte-vectorization utilities added by that PR.
Suggested reviewers: bernhardmgruber, jacobfaib
Priority:
Change: Feature
Merge Risk: 
3f927
-
Update the loader documentation and confirm the outstanding cache-load, conversion, and alignment behavior before merging.
Comment @coderabbitai help to get the list of available commands.
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments: 1
cub/cub/block/block_load.cuh: 1
227-233: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Update the
LoadDirectBlockedVectorizeddocumentation and theBLOCK_LOAD_VECTORIZEandWARP_LOAD_VECTORIZEenum documentation. They no longer match the new behavior.
cuda::__as_bytes_vectorized_t<T, ItemsPerThread>packs oddItemsPerThreadinto word-sized vectors. An odd count no longer forces the direct fallback.- Any
Twithalignof(T) <= __max_vector_accessis vectorizable. This is not limited to built-in primitives and CUDA vector types.- The required alignment is
alignof(vector_t)forvector_t = cuda::__as_bytes_vectorized_t<T, ItemsPerThread>. The text "quad-item aligned" and "quadword-aligned" is no longer correct.Stale text is in
block_load.cuhLines 227-233 and 671-678, and inwarp_load.cuhLines 86-90.As per coding guidelines: "The
@brief/@param/@returndescription must accurately reflect the current functionality of the function."Source: Coding guidelines
ℹ️ Review info
⚙️ Run configuration
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: as&
#95;vectorized_bytes - File: .coderabbit.yaml
- Commit: 2640363
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
a9b73ab7-9079-4946-a28c-89b440b8d153
📒 Files selected for processing: 2
cub/cub/block/block_load.cuhcub/cub/warp/warp_load.cuh
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 5 remain after this review.
| //! @brief Internal implementation of LoadDirectBlockedVectorized for input aligned to the vector type | ||
| //! | ||
| //! @param[in] linear_tid | ||
| //! A suitable 1D thread-identifier for the calling thread (e.g., `(threadIdx.y * blockDim.x) + linear_tid` for 2D | ||
| //! thread blocks) | ||
| //! | ||
| //! @param[in] block_src_ptr | ||
| //! Input pointer for loading from | ||
| //! Input pointer for loading from, aligned to ``cuda::__as_bytes_vectorized_t<T, ItemsPerThread>`` | ||
| //! | ||
| //! @param[out] dst_items | ||
| //! Destination to load data into | ||
| template <CacheLoadModifier MODIFIER, typename T, int ItemsPerThread> | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void | ||
| InternalLoadDirectBlockedVectorized(int linear_tid, const T* block_src_ptr, T (&dst_items)[ItemsPerThread]) | ||
| { | ||
| // Find biggest memory access word that T is a whole multiple of | ||
| using device_word_t = typename UnitWord<T>::DeviceWord; | ||
| _CCCL_DIAG_PUSH | ||
| # if _CCCL_COMPILER(CLANG, >=, 10) | ||
| _CCCL_DIAG_SUPPRESS_CLANG("-Wsizeof-array-div") | ||
| # endif // _CCCL_COMPILER(CLANG, >=, 10) | ||
| // NOLINTNEXTLINE(bugprone-sizeof-expression) | ||
| constexpr int total_words = static_cast<int>(sizeof(dst_items) / sizeof(device_word_t)); | ||
| _CCCL_DIAG_POP | ||
| constexpr int vector_size = (total_words % 4 == 0) ? 4 : (total_words % 2 == 0) ? 2 : 1; | ||
| constexpr int vectors_per_thread = total_words / vector_size; | ||
| // All items of a thread are loaded with the widest vectors that evenly divide them | ||
| using vector_t = ::cuda::__as_bytes_vectorized_t<T, ItemsPerThread>; | ||
| _CCCL_ASSERT(::cuda::std::is_sufficiently_aligned<alignof(vector_t)>(block_src_ptr), | ||
| "block_src_ptr must be aligned to the vector type"); | ||
|
|
||
| // Load into an array of vectors in thread-blocked order | ||
| using vector_t = typename CubVector<device_word_t, vector_size>::Type; | ||
| // Load in thread-blocked order | ||
| const vector_t vec_items = ThreadLoad<MODIFIER>(reinterpret_cast<const vector_t*>(block_src_ptr) + linear_tid); | ||
|
|
||
| // Add the alignment check to ensure the vectorized loading can proceed. | ||
| if (::cuda::std::is_sufficiently_aligned<alignof(vector_t)>(block_src_ptr)) | ||
| { | ||
| vector_t vec_items[vectors_per_thread]; | ||
| // Load into an array of vectors in thread-blocked order | ||
| const vector_t* vec_ptr = reinterpret_cast<const vector_t*>(block_src_ptr) + linear_tid * vectors_per_thread; | ||
|
|
||
| _CCCL_PRAGMA_UNROLL_FULL() | ||
| for (int i = 0; i < vectors_per_thread; i++) | ||
| { | ||
| vec_items[i] = ThreadLoad<MODIFIER>(vec_ptr + i); | ||
| } | ||
|
|
||
| // Copy to destination | ||
| _CCCL_PRAGMA_UNROLL_FULL() | ||
| for (int i = 0; i < ItemsPerThread; i++) | ||
| { | ||
| dst_items[i] = *(reinterpret_cast<T*>(vec_items) + i); | ||
| } | ||
| } | ||
| else | ||
| // Copy to destination | ||
| _CCCL_PRAGMA_UNROLL_FULL() | ||
| for (int i = 0; i < ItemsPerThread; i++) | ||
| { | ||
| LoadDirectBlocked(linear_tid, block_src_ptr, dst_items); | ||
| dst_items[i] = vec_items.template __alias<T>(i); |
There was a problem hiding this comment.
🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win
🔎 Supported by static analysis
-
🏁 Script executed:
sed -n '180,290p' cub/cub/block/block_load.cuh sed -n '325,360p' cub/cub/warp/warp_load.cuh rg -n 'ThreadLoad|__as_bytes_vectorized_t' cub/cub/thread libcudacxx/include/cuda | head -110
Repository: NVIDIA/cccl
Length of output: 11358
🏁 Script executed:
set -o pipefail printf '%s\n' '--- thread_load definitions ---' nl -ba cub/cub/thread/thread_load.cuh | sed -n '130,390p' printf '%s\n' '--- byte-vector type definition ---' nl -ba libcudacxx/include/cuda/__memory/as_bytes_vectorized.h | sed -n '1,180p' printf '%s\n' '--- current cache-modified BlockLoad call sites ---' rg -n -F --glob '*.cuh' -e 'LoadDirectBlockedVectorized<' -e 'CacheModifiedInputIterator' cub/cub/block cub/cub/warp printf '%s\n' '--- current BlockLoad load dispatch ---' rg -n -F --glob 'block_load.cuh' -e 'LoadDirectBlockedVectorized' -e 'CacheModifiedInputIterator' -e 'template <CacheLoadModifier' cub/cub/block/block_load.cuh printf '%s\n' '--- merge-base relevant implementation ---' git show 264036356b1270d075ea4842b54af06cb41644e6:cub/cub/block/block_load.cuh | nl -ba | sed -n '155,285p' printf '%s\n' '--- current diff for relevant files ---' git diff --unified=35 264036356b1270d075ea4842b54af06cb41644e6 1fec506b496a3859cfab50521aa41fedd49dcd8d -- cub/cub/block/block_load.cuh cub/cub/warp/warp_load.cuh
Repository: NVIDIA/cccl
Length of output: 41873
🏁 Script executed:
nl -ba cub/cub/thread/thread_load.cuh | sed -n '270,370p' nl -ba libcudacxx/include/cuda/__memory/as_bytes_vectorized.h | sed -n '1,160p' nl -ba cub/cub/block/block_load.cuh | sed -n '850,1050p' nl -ba cub/cub/warp/warp_load.cuh | sed -n '330,385p' git diff --unified=25 264036356b1270d075ea4842b54af06cb41644e6 1fec506b496a3859cfab50521aa41fedd49dcd8d -- cub/cub/block/block_load.cuh cub/cub/warp/warp_load.cuh
Repository: NVIDIA/cccl
Length of output: 38156
🏁 Script executed:
printf '%s\n' '--- UnitWord implementation ---' nl -ba cub/cub/detail/unit_word.cuh | sed -n '1,260p' printf '%s\n' '--- related UnitWord definitions/usages ---' rg -n -F --glob '*.cuh' -e 'UnitWord<' -e 'DeviceWord' -e 'VolatileWord' cub/cub | head -180 printf '%s\n' '--- current vectorized helper and dispatch ---' nl -ba cub/cub/block/block_load.cuh | sed -n '180,285p' nl -ba cub/cub/thread/thread_load.cuh | sed -n '338,382p'
Repository: NVIDIA/cccl
Length of output: 16875
🏁 Script executed:
nl -ba cub/cub/detail/unit_word.cuh | sed -n '1,240p'
Repository: NVIDIA/cccl
Length of output: 5325
important: Align the temporary cache-load buffer to T.
On CTK 13+, a 32-byte byte-vector type can reach the nondefault ThreadLoad path. That path stores words in an array aligned only for DeviceWord, then reinterprets it as the 32-byte-aligned vector type. This can perform a misaligned access before __alias copies the items.
Suggested fix
- LoadWord words[LOAD_MULTIPLE];
+ alignas(T) LoadWord words[LOAD_MULTIPLE];
This comment has been minimized.
This comment has been minimized.
| // Load into an array of vectors in thread-blocked order | ||
| using vector_t = typename CubVector<device_word_t, vector_size>::Type; | ||
| // Load in thread-blocked order | ||
| const vector_t vec_items = ThreadLoad<MODIFIER>(reinterpret_cast<const vector_t*>(block_src_ptr) + linear_tid); |
There was a problem hiding this comment.
Note: This does a load of the whole vector_t struct.
This gives the compiler the ability to structure that load as it deems necessary. In all the cases we have internally, this meant that we used the widest loads possible
This comment has been minimized.
This comment has been minimized.
1fec506 to
e4c04b1
Compare
e4c04b1 to
95a26ff
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
-
-
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: as&
#95;vectorized_bytes - File: .coderabbit.yaml
- Commit: f82b645
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
434f3464-425b-4b4b-bb23-e88620b50994
-
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; 10 remain after this review.
-
| for (int i = 0; i < ItemsPerThread; i++) | ||
| { | ||
| LoadDirectBlocked(linear_tid, block_src_ptr, dst_items); | ||
| dst_items[i] = vec_items.template __alias<InputT>(i); |
There was a problem hiding this comment.
important: Convert user-defined inputs from the source object.
When InputT differs from T, this conversion runs on vec_items, not on the input object. A trivially copyable InputT can still define an operator T() that increments a mutable field or depends on its address. The new vectorized path changes that conversion’s result or side effects; the previous dispatch used direct source-object loads. Keep direct loading for user-defined conversions. (raw.githubusercontent.com)
| InternalLoadDirectBlockedVectorized<LOAD_DEFAULT>(linear_tid, block_src_ptr, dst_items); | ||
| if constexpr (!::cuda::__is_vectorizable<InputT>()) | ||
| { | ||
| LoadDirectBlocked(linear_tid, block_src_ptr, dst_items); |
There was a problem hiding this comment.
important: Preserve cache modifiers in direct fallbacks.
If a cache-modified input has a different destination type and misses vectorization, these branches load through block_src_ptr instead of the iterator. For example, a misaligned CacheModifiedInputIterator<LOAD_VOLATILE, int> loaded into long long no longer makes volatile loads. The previous dispatch dereferenced that iterator for differing types. Pass the modifier to each direct load, or retain the iterator for fallback loads. (raw.githubusercontent.com)
Also applies to: 273-273
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
95a26ff to
5708f7a
Compare
5708f7a to
85ef93d
Compare
There was a problem hiding this comment.
Actionable comments posted: 1
-
-
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: as&
#95;vectorized_bytes - File: .coderabbit.yaml
- Commit: ae65cdf
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
f9081d4b-1c8d-4410-888b-eca58cc408cc
-
cub/cub/block/block_load.cuh
Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 9 remain after this review.
-
| { | ||
| using vector_t = ::cuda::__as_bytes_vectorized_t<InputT, ItemsPerThread>; | ||
| // Add the alignment check to ensure the vectorized loading can proceed. | ||
| if (::cuda::std::is_sufficiently_aligned<alignof(vector_t)>(block_src_ptr)) |
There was a problem hiding this comment.
🔎 Supported by static analysis
-
Script executed:
# Head: dispatch around the reviewed line and the vector_t definition. nl -ba cub/cub/block/block_load.cuh | sed -n '225,285p' rg -n -F -e 'vector_t' -e 'LoadDirectBlockedVectorized' -- cub/cub/block/block_load.cuh || test "$?" -eq 1 # Head: alignment declarations in the vectorized-bytes header. rg -n -e 'align' -e 'vector' -- libcudacxx/include/cuda/__memory/as_bytes_vectorized.h || test "$?" -eq 1 # Base: the previous LoadDirectBlockedVectorized implementation and its vector width. git show ae65cdfc94dc8683660a69f20920efaa4e547b80:cub/cub/block/block_load.cuh | nl -ba | rg -n -e 'LoadDirectBlockedVectorized' -e 'vector' -e 'is_sufficiently_aligned' -e 'uint4' -e 'int4' || test "$?" -eq 1
Repository: NVIDIA/cccl
Length of output: 8745
important: Keep a 16-byte vector path for 16-byte-aligned input.
For int with ItemsPerThread=8, vector_t is 32-byte aligned when the vector access limit is 32 bytes, as on CUDA Toolkit 13. The check at block_load.cuh:267 then rejects any input that is only 16-byte aligned. The call falls back to LoadDirectBlocked, which issues scalar loads.
The base implementation handled this case. Its InternalLoadDirectBlockedVectorized used a 16-byte vector_t with two vectors per thread, so 16-byte-aligned input got two 16-byte loads per thread. The new dispatch removes that path, so this input can regress to scalar loads.
Add a 16-byte vector path before the scalar fallback. Use it when the input is 16-byte aligned but not 32-byte aligned.
Source: Path instructions
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
This comment has been minimized.
85ef93d to
3f9275d
Compare
There was a problem hiding this comment.
Caution
Some comments are outside the diff and can’t be posted inline due to GitHub limitations.
block_load.cuh:245-249
cub/cub/block/block_load.cuh:245-249
📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick winimportant: The
LoadDirectBlockedVectorizeddocumentation no longer matches the implementation.
- Line 243 says the offset must be quad-item aligned. The code now requires
alignofof the vector type and falls back at run time.- Line 247 says an odd
ItemsPerThreadprevents vectorization. The code now builds the vector fromItemsPerThreadorVecSizeand has no parity check.- Lines 248-249 describe
Tas the eligibility criterion. Eligibility now depends onInputTthroughblock_load_can_vectorize_v: it must be primitive or trivially copyable, non-volatile, and read through a contiguous pointer.The
BLOCK_LOAD_VECTORIZEenumerator documentation at lines 715-720 has the same stale conditions. Update both.As per coding guidelines: "The
@brief/@param/@returndescription must accurately reflect the current functionality of the function."Source: Coding guidelines
-
-
- Source mode: Auto
- Repository: NVIDIA/cccl
- Branch: vectorized&
#95;block_load - File: .coderabbit.yaml
- Commit: 3f9275d
- Layers: Repository YAML
- Review profile: CHILL
- Plan: Enterprise
- Run ID:
7e75d063-e57f-46a2-ab58-0a9c4a16e12a
-
cub/cub/agent/agent_topk.cuhcub/cub/block/block_load.cuhcub/cub/device/dispatch/dispatch_topk.cuhcub/cub/device/dispatch/tuning/tuning_topk.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.
`InternalLoadDirectBlockedVectorized` now loads all items of a thread with the widest vectors that evenly divide them and requires the input to be aligned to `cuda::__as_bytes_vectorized_t<InputT, ItemsPerThread>`. The alignment check and the fallback to `LoadDirectBlocked` move into `LoadDirectBlockedVectorized`, which also falls back for types that cannot be vectorized, takes a pointer to const and an optional cache load modifier. The vectors are built from the value type `InputT` of the input and each item is converted to `T` when it is copied out of them. Therefore, `BLOCK_LOAD_VECTORIZE` no longer requires the value type of the iterator to be `T`. It also no longer checks contiguous iterators for volatile, because `cuda::__is_vectorizable` rejects volatile types and `LoadDirectBlockedVectorized` falls back to `LoadDirectBlocked` for them. `WarpLoad` called `InternalLoadDirectBlockedVectorized` directly and relied on its fallback for unaligned input, so it switches to `LoadDirectBlockedVectorized` in the same commit.
3f9275d to
ea9b305
Compare
🔬 Benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
😬 CI Workflow Results🟥 Finished in 2h 42m: Pass: 98%/210 | Total: 6d 19h | Max: 2h 34m | Hits: 59%/567650See results here. AI failure analysis1. C Parallel v2 transform segfault in integral case · 1 jobExplanation: The Windows HostJIT transform executable crashes during the first generated integral test case. Source line `test_util.h:735` is the host-to-device `cudaMemcpy`; because GitHub provided no crash stack, it is unclear whether the zero-item case itself triggers the runtime fault or earlier corruption reaches this call, and the relationship to the PR's low-level CUB load changes is not yet established. Evidence: Copy this prompt into a coding agentJobs: 2. CUB SASS changes lack required approval · 1 jobExplanation: The approval job found no SASS-approval artifact for the current commit. This is an independent policy gate, not a derivative test failure; the PR changes vectorized loads and tuning, so generated instruction changes are plausible, but the log does not establish whether they are intended. Evidence: Copy this prompt into a coding agentJobs: 3. Python host benchmark incompatible with pytest 9.2 · 3 jobsExplanation: All three jobs resolve the same unbounded dependency pair, and `pytest-benchmark` 5.3.0 accesses pytest's removed private `Function._nodeid` attribute while creating every benchmark fixture. The benchmark bodies never execute; this is dependency drift rather than a CUDA backend or platform-specific failure. Evidence: Copy this prompt into a coding agentJobs: |
This uses the new machinery for BlockLoad and WarpLoad
Stack created with GitHub Stacks CLI • Give Feedback 💬