[cudax] Implement cudax::level_synchronizer#9871
Open
davebayer wants to merge 1 commit into
Open
Conversation
Contributor
📝 WalkthroughSummary by CodeRabbit
WalkthroughAdds ChangesLevel Synchronizer
Group Synchronization Delegation
Synchronizer Interface Validation
Assessment Against Linked Issues
Comment |
![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: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: 51b715af-9d50-4229-a812-ded9b32ee448
📒 Files selected for processing (12)
cudax/include/cuda/experimental/__group/fwd.cuhcudax/include/cuda/experimental/__group/group.cuhcudax/include/cuda/experimental/__group/synchronizer/barrier_synchronizer.cuhcudax/include/cuda/experimental/__group/synchronizer/lane_synchronizer.cuhcudax/include/cuda/experimental/__group/synchronizer/level_synchronizer.cuhcudax/include/cuda/experimental/__group/this_group.cuhcudax/include/cuda/experimental/group.cuhcudax/test/CMakeLists.txtcudax/test/group/synchronizer/barrier_synchronizer.cucudax/test/group/synchronizer/lane_synchronizer.cucudax/test/group/synchronizer/level_synchronizer.cucudax/test/group/this_group.cu
Contributor
🥳 CI Workflow Results🟩 Finished in 48m 18s: Pass: 100%/57 | Total: 15h 18m | Max: 48m 15s | Hits: 26%/128114See results here. |
Jacobfaib
reviewed
Jul 16, 2026
Comment on lines
+44
to
+48
| [[nodiscard]] _CCCL_DEVICE_API auto | ||
| make_instance(const _Unit&, const _ParentGroup&, const _MappingResult&) const noexcept | ||
| { | ||
| return __synchronizer_instance<typename _ParentGroup::level_type>{}; | ||
| } |
Contributor
There was a problem hiding this comment.
Suggested change
| [[nodiscard]] _CCCL_DEVICE_API auto | |
| make_instance(const _Unit&, const _ParentGroup&, const _MappingResult&) const noexcept | |
| { | |
| return __synchronizer_instance<typename _ParentGroup::level_type>{}; | |
| } | |
| [[nodiscard]] _CCCL_DEVICE_API constexpr __synchronizer_instance<typename _ParentGroup::level_type> | |
| make_instance(const _Unit&, const _ParentGroup&, const _MappingResult&) const noexcept | |
| { | |
| return {}; | |
| } |
(the constexpr is important, but feel free to reject the rest)
| const auto __barrier_ptr = &__grid_workspace_ptr->__barrier_; | ||
|
|
||
| // Synchronize the block before synchronizing with the other blocks. | ||
| ::cuda::experimental::__block_sync<_Aligned>(); |
Contributor
There was a problem hiding this comment.
Hmmm you've done all this typing implementing the various level synchronizers, but you still just call the bare helper function here (and above in the cluster_level impl). Wouldn't it be cleaner to reuse those types here? Would also help to dogfood the API
Comment on lines
+199
to
+207
| unsigned __curr_barrier_value; | ||
| do | ||
| { | ||
| # if _CCCL_HAS_NV_ATOMIC_BUILTINS() | ||
| __nv_atomic_load(__barrier_ptr, &__curr_barrier_value, __NV_ATOMIC_ACQUIRE, __NV_THREAD_SCOPE_DEVICE); | ||
| # else // ^^^ _CCCL_HAS_NV_ATOMIC_BUILTINS() ^^^ / vvv !_CCCL_HAS_NV_ATOMIC_BUILTINS() vvv | ||
| asm volatile("ld.acquire.gpu.u32 %0, [%1];" : "=r"(__curr_barrier_value) : "l"(__barrier_ptr) : "memory"); | ||
| # endif // ^^^ !_CCCL_HAS_NV_ATOMIC_BUILTINS() ^^^ | ||
| } while (static_cast<int>(__old_barrier_value) < 0 == static_cast<int>(__curr_barrier_value) < 0); |
Contributor
There was a problem hiding this comment.
Surely we can reuse cuda::std::atomic here no? Or cuda::std::barrier?
| : __hier_{::cuda::__unpack_hierarchy_if_needed(__hier_like)} | ||
| {} | ||
|
|
||
| [[nodiscard]] _CCCL_DEVICE_API const hierarchy_type& hierarchy() const noexcept |
Contributor
There was a problem hiding this comment.
Suggested change
| [[nodiscard]] _CCCL_DEVICE_API const hierarchy_type& hierarchy() const noexcept | |
| [[nodiscard]] _CCCL_DEVICE_API constexpr const hierarchy_type& hierarchy() const noexcept |
(and below as well)
| # if _CCCL_HAS_COOPERATIVE_GROUPS() | ||
| # if _CCCL_HAS_COOPERATIVE_GROUPS() | ||
| template <class _Parent> | ||
| _CCCL_DEVICE_API this_thread(const ::cooperative_groups::thread_block_tile<1, _Parent>&) noexcept |
Contributor
There was a problem hiding this comment.
This fella could be constexpr as well (should do a pass over these APIs in a follow up PR)
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
2 participants
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.
Fixes #9870