Repository navigation
Conversation
Contributor
|
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. |
This comment was marked as resolved.
This comment was marked as resolved.
Contributor
🔬 Benchmark SASS comparisonHow to request a benchmark run
Targets with a SASS change
|
Contributor
Author
Perf results on B200 for variable segment sizes indexed TopK w/o re-tuningNote that the regressions for small sizes are in an area that is handled by the non-cluster backend. According to the robust comparison the data in that area is not to be trusted anyway. Legacy Comparison['/home/pgrossebley/scratch/b200_134_seg_topk_var_ind_cluster-10-24_tie-large_ref.json', '/home/pgrossebley/scratch/b200_134_seg_topk_var_ind_cluster-10-24_tie-large_010.json'] decode_style_variable_topk_indexed[0] NVIDIA B200Summary
The robust comparison table did not fit into the same comment due to the number of rows per table. See next comment. |
Contributor
Author
Robust Comparisondecode_style_variable_topk_indexed[0] NVIDIA B200Summary
|
Contributor
Author
Perf results on B200 for variable segment sizes keys-only TopK w/o re-tuningLegacy Comparison['/home/pgrossebley/scratch/b200_134_seg_topk_var_keys_cluster-10-24_tie-large_ref.json', '/home/pgrossebley/scratch/b200_134_seg_topk_var_keys_cluster-10-24_tie-large_010.json'] decode_style_variable_topk_keys[0] NVIDIA B200Summary
|
Contributor
Author
Robust Comparisondecode_style_variable_topk_keys[0] NVIDIA B200Summary
Details |
Contributor
🥳 CI Workflow Results🟩 Finished in 3h 19m: Pass: 100%/210 | Total: 9d 20h | Max: 3h 19m | Hits: 22%/1254980See results here. |
miscco
reviewed
Oct 26, 2026
| size_state size; | ||
| pass_result result; | ||
| }; | ||
| static_assert(sizeof(state_t) == 16, "the published cluster state must be exactly one 16-byte st.async payload"); |
Contributor
There was a problem hiding this comment.
I believe this is superfluous if we manually just added the alignas
Contributor
Author
There was a problem hiding this comment.
Is it? Maybe I'm missing sth, but aggregate types that are bigger than their alignment are not unusual. Did you misread the assert to be on alignof()?
Comment on lines
+850
to
+865
| // mbarrier-path `add_remote_prefix`, completing on the target's `prefix_mbar` (see `exchange_prefixes`). | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void add_remote_prefix_async( | ||
| [[maybe_unused]] unsigned target_rank, | ||
| [[maybe_unused]] offset_t num_local_selected_for_scan, | ||
| [[maybe_unused]] offset_t num_local_tied_for_scan) | ||
| { | ||
| #if __cccl_ptx_isa >= 860 | ||
| _TempStorage* const remote = map_to_rank(&temp_storage, target_rank); | ||
| ::cuda::ptx::red_async( | ||
| ::cuda::ptx::op_add, &remote->selected_offset_counter, num_local_selected_for_scan, &remote->prefix_mbar); | ||
| ::cuda::ptx::red_async( | ||
| ::cuda::ptx::op_add, &remote->tie_offset_counter, num_local_tied_for_scan, &remote->prefix_mbar); | ||
| #else // __cccl_ptx_isa < 860 | ||
| _CCCL_UNREACHABLE(); | ||
| #endif // __cccl_ptx_isa >= 860 | ||
| } |
Contributor
There was a problem hiding this comment.
We might want to conditionally compile the whole function and not just the internal.
Contributor
Author
There was a problem hiding this comment.
I guess that should help compile time? On the other hand I do like having the guards directly next to the PTX in question.
Contributor
Author
There was a problem hiding this comment.
I might just inline some of these new helpers. Small functions like these seem nice at first sight, but when reviewing complicated kernels they make you jump around in the source a lot and potentially add a lot of boilerplate like these guards and the function signatures. Often times they are only used in one place anyway.
| ::cuda::ptx::fence_sync_restrict(::cuda::ptx::sem_acquire, ::cuda::ptx::space_cluster, ::cuda::ptx::scope_cluster); | ||
| #else // __cccl_ptx_isa < 860 | ||
| _CCCL_UNREACHABLE(); | ||
| #endif // __cccl_ptx_isa >= 860 |
Contributor
There was a problem hiding this comment.
A lot of this feels like an anti pattern. Can you please check if we cannot just rely on use_mbarrier_sync with an if constexpr and NV_IF_TARGET to conditionally compile out all those functions rather than having a lot of unreachable functions?
I
Contributor
Author
There was a problem hiding this comment.
AFAIK the issue is that these functions are not templates, so even when "guarded" by if constexpr on the caller side they cause build failures. But I will take another look.
| // plus the remote relaxed arrival (which also arms the receiver for the payload) is what orders this pass's `hist` | ||
| // accesses (BlockScan reads, reset) before the receiver's next fold. The caller's block barrier puts every thread's | ||
| // accesses into the issuing lanes' history. | ||
| _CCCL_DEVICE _CCCL_FORCEINLINE void publish_pass_state() |
Contributor
There was a problem hiding this comment.
Same here, this is fully guarded by use_mbarrier_sync so we can just conditionally define the whole function
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.
Description
Certain features used on the new
mbarrier-based fast path require PTX ISA 8.6 (has no corresponding public CTK, CTK 12.8 gives PTX ISA 8.7). Therefore the originalbarrier.cluster-based path is currently preserved for older CTKs. For CTK 12.1-12.6 we could instead use a slowermbarrier-based path that has fences doing L2 roundtrips but might still profit from more fine-grained synchronization. As long as we keep support for CTK 12.0 we need to keep thebarrier.cluster-path though since CTK 12.0 does not yet havered.async/st.asyncwhich we need in thembarrier-based path. So for now I plan to leave it in the current state as long as there is no request for (maybe) better CTK 12.1-12.6 performance. Especially since we have not tuned yet for Hopper in the first place and this additional work would only be help on Hopper.closes #9994
Checklist