[cudax] Rework `cudax::coop::barrier_synchronizer` by davebayer · Pull Request #12229 · NVIDIA/cccl

davebayer · 2026-10-23T18:32:27Z

This PR rewords cudax::coop::barrier_synchronizer to take a range of barriers instead of a contiguous array.

coderabbitai · 2026-10-23T18:39:51Z

📝 Summary

Summary by CodeRabbit

New Features
- The barrier synchronizer accepts viewable ranges of CUDA barriers, including reversed ranges, and assigns barriers to groups by position.
- Both sized and unsized ranges are supported, with checks that enough barriers are available for the groups.
API Changes
- The synchronizer now uses a range view as its template parameter instead of a barrier type and span extent. The barrier_type alias and barriers() accessor are no longer available.
- make_instance is no longer const.
Documentation
- Updated usage guidance to describe barrier ordering, count, accessibility, and scope requirements.

barrier_synchronizer now accepts viewable ranges of CUDA barriers and stores their views::all view. Instance creation selects barriers by group rank. Tests and documentation cover range construction and selection.

Changes: Barrier synchronizer range support

Layer / File(s)	Summary
Range-view API `cudax/include/cuda/experimental/coop/__group/fwd.cuh`, `cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh`, `docs/cudax/coop/group/synchronizer/barrier_synchronizer.rst`	The template parameter changes to a range view. The constructor stores `views::all` of a viewable range, and the deduction guide accepts ranges whose value type is a CUDA barrier. The `barrier_type` alias and `barriers()` accessor are removed.
Barrier selection by group rank `cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh`, `cudax/test/coop/group/synchronizer/barrier_synchronizer.cu`, `cudax/test/coop/group/generic_group.cu`, `docs/cudax/coop/group/synchronizer/barrier_synchronizer.rst`	Instance creation selects the barrier at `group_rank()`. Sized ranges check that the group count does not exceed the range size. Unsized ranges check that iteration reaches the requested rank. Tests and examples cover range construction, including a reverse view. The generic group tests use non-const synchronizers.

Suggested reviewers: jacobfaib, elstehle

Priority: ➖ Normal

Change: Feature

Merge Risk: 🔵 Low · up to 39ff2

The range API has a conditional compile failure and the tests need explicit template arguments before merge.

- [ ] Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

_{Comment @coderabbitai help to get the list of available commands.}

✨ Finishing Touches: 💡 1

🛠️ Fix failing CI checks: 💡

Commit to this branch
Create a new PR

coderabbitai

Actionable comments posted: 2

ℹ️ Review info

⚙️ Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 4811a1c
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: 5edcfec5-b438-4a5a-b137-c63cb20648ae

📥 Commits

Reviewing files that changed from the base of the PR and between 4811a1c and e5dd794.

📒 Files selected for processing: 4

cudax/include/cuda/experimental/coop/__group/fwd.cuh
cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh
cudax/test/coop/group/synchronizer/barrier_synchronizer.cu
docs/cudax/coop/group/synchronizer/barrier_synchronizer.rst

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 11 remain after this review.

Jacobfaib · 2026-10-23T18:43:36Z

+  _CCCL_REQUIRES(::cuda::std::is_same_v<_RangeView, ::cuda::std::ranges::views::all_t<_Range>>)
+  _CCCL_DEVICE_API barrier_synchronizer(_Range&& __barriers) noexcept(
+    noexcept(::cuda::std::ranges::views::all(::cuda::std::forward<_Range>(__barriers))))
+      : __barriers_{::cuda::std::ranges::views::all(::cuda::std::forward<_Range>(__barriers))}


If it's already all_t, why do you need to rewrap it in all()? Or am I misreading this?

No, it's a viewable range that I store as all_t<R> range

Jacobfaib · 2026-10-23T18:48:19Z

+_CCCL_REQUIRES(
+  ::cuda::std::ranges::viewable_range<_Range> _CCCL_AND __is_cuda_barrier_v<::cuda::std::ranges::range_value_t<_Range>>)
+_CCCL_DEDUCTION_GUIDE_ATTRIBUTES barrier_synchronizer(_Range&&)
+  -> barrier_synchronizer<::cuda::std::ranges::views::all_t<_Range>>;


Why all_t here? Why not just _Range?

Because it handles for me all of the problems that come with taking a view over a range. I discussed this in the past with @miscco, it's the most suitable thing

Jacobfaib · 2026-10-23T18:50:51Z

    if (__mapping_result.is_valid())
    {
-      __group_barrier_ptr = __barriers_.data() + __mapping_result.group_rank();
+      if constexpr (::cuda::std::ranges::sized_range<_RangeView>)


For your own sanity, I would suggest requiring sized ranges. I guess it might preclude filter views but not having constant-time access to size of the range also means that the range could be forward-pass only, or have some other weird quirks.

Jacobfaib · 2026-10-23T18:57:58Z

+  static_assert(__is_cuda_barrier_v<::cuda::std::ranges::range_value_t<_RangeView>>,
+                "_RangeView's value type must be cv-unqualified cuda::barrier type");


More generally, I would impose far more restrictions on the range here as well. It seems like at the very least you want:

RangeView satisfies ranges::forward_range (you seemingly want to be able to iterate a few times over the range)

You probably also want ranges::borrowed_range since you'll be storing the view and so can't e.g. accept a temporary or initializer list.

You do not actually care about the range_value_t. What you care about is range_reference_t because you do not store the value, you only addressof(*it).

I would also say we need both forward_range and borrowed_range

coderabbitai

Actionable comments posted: 1

🧹 Nitpick comments: 1

cudax/test/coop/group/synchronizer/barrier_synchronizer.cu: 1

48-55: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Cover the unsized range and barrier selection paths.

The make_instance test uses only a sized array. Add a non-sized forward range to exercise the iterator path. Also verify that different group_rank() values select different barriers. Existing related tests use barrier arrays, while the unsized-range selection test covers only interwarp_synchronizer.

ℹ️ Review info

⚙️ Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 4811a1c
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: 1fbdd949-9d3a-45e4-bfe0-aecc1da75eb2

📥 Commits

Reviewing files that changed from the base of the PR and between e5dd794 and f9a5ef6.

📒 Files selected for processing: 2

cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh
cudax/test/coop/group/synchronizer/barrier_synchronizer.cu

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 7 remain after this review.

coderabbitai

Actionable comments posted: 1

ℹ️ Review info

⚙️ Run configuration

Source mode: Auto
Repository: NVIDIA/cccl
Branch: main
File: .coderabbit.yaml
Commit: 4811a1c
Layers: Repository YAML
Review profile: CHILL
Plan: Enterprise
Run ID: 8b5bda43-2dfa-427c-9b71-750ca952100e

📥 Commits

Reviewing files that changed from the base of the PR and between f9a5ef6 and 39ff271.

📒 Files selected for processing: 1

cudax/test/coop/group/generic_group.cu

Included review availability: This review used your included allowance. Your plan provides up to 12 included reviews per hour; 6 remain after this review.

coderabbitai · 2026-10-23T19:23:34Z


    const cudax::coop::group_by mapping{cuda::std::cw<N>};
-    const cudax::coop::barrier_synchronizer synchronizer{barriers};
+    cudax::coop::barrier_synchronizer synchronizer{barriers};


📐 Maintainability & Code Quality | 🟡 Minor | ⚡ Quick win

important: Use explicit template arguments for barrier_synchronizer.

Both declarations use barrier_synchronizer synchronizer{barriers} without template arguments. The coding guideline forbids deduction-guide-dependent initialization. Spell out the _RangeView template argument in both tests.

Also applies to: 125-125

Source: Coding guidelines

pciolkosz · 2026-10-23T22:57:03Z

I kinda fail to see why this is useful. Is there some specific case that inspired this? In general I think usage of explicit synchronizers will be rare and this seems even more obscure

davebayer · 2026-10-26T06:09:37Z

I kinda fail to see why this is useful. Is there some specific case that inspired this? In general I think usage of explicit synchronizers will be rare and this seems even more obscure

I'd like all of the groups APIs to be somewhat uniform. All mappings that take multiple values are implemented to accept a C++ range. I think it would make sense to do the same thing for this case as well

davebayer requested review from a team as code owners October 23, 2026 18:32

davebayer requested a review from fbusato October 23, 2026 18:32

github-project-automation Bot added this to CCCL Oct 23, 2026

github-project-automation Bot moved this to Todo in CCCL Oct 23, 2026

cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Oct 23, 2026

coderabbitai Bot reviewed Oct 23, 2026

View reviewed changes

Comment thread cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh

Comment thread cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh

coderabbitai Bot mentioned this pull request Oct 23, 2026

[cudax] Rework cudax::coop::group_as mapping #12223

Merged

Jacobfaib reviewed Oct 23, 2026

View reviewed changes

[cudax] Rework cudax::coop::barrier_synchronizer

f9a5ef6

davebayer force-pushed the rework_barrier_synchronizer branch from e5dd794 to f9a5ef6 Compare October 23, 2026 19:03

coderabbitai Bot reviewed Oct 23, 2026

View reviewed changes

Comment thread cudax/include/cuda/experimental/coop/__group/synchronizer/barrier_synchronizer.cuh

fix

39ff271

coderabbitai Bot reviewed Oct 23, 2026

View reviewed changes

		static_assert(__is_cuda_barrier_v<::cuda::std::ranges::range_value_t<_RangeView>>,
		"_RangeView's value type must be cv-unqualified cuda::barrier type");

Conversation

davebayer commented Oct 23, 2026

Uh oh!

coderabbitai Bot commented Oct 23, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary by CodeRabbit

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Jacobfaib Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

davebayer Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

davebayer Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

miscco Oct 26, 2026

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

coderabbitai Bot left a comment

Choose a reason for hiding this comment

Uh oh!

coderabbitai Bot Oct 23, 2026

Choose a reason for hiding this comment

Uh oh!

pciolkosz commented Oct 23, 2026

Uh oh!

davebayer commented Oct 26, 2026

Uh oh!

This branch has not been deployed

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

coderabbitai Bot commented Oct 23, 2026 •

edited

Loading