Skip to content

[cudax] Expose implict_hierarchy[_1d] - #10911

Merged
davebayer merged 1 commit into
NVIDIA:mainfrom
davebayer:native_hierarchies
Aug 20, 2026
Merged

[cudax] Expose implict_hierarchy[_1d]#10911
davebayer merged 1 commit into
NVIDIA:mainfrom
davebayer:native_hierarchies

Conversation

@davebayer

Copy link
Copy Markdown
Contributor

This PR makes the implicit_hierarchy public and adds the implicit_hierarchy_1d function for generating dynamic 1D hierarchy. This will be useful for users that don't use cuda::launch to use the groups and need to construct the hierarchy themselves on device.

I've also removed the default constructor from this_meow groups. User should be aware he is using the dynamic hierarchy and use implicit_hierarchy_1d whenever possible to optimize the queries.

@davebayer
davebayer requested a review from a team as a code owner August 20, 2026 08:28
@davebayer
davebayer requested a review from andralex August 20, 2026 08:28
@github-project-automation github-project-automation Bot moved this to Todo in CCCL Aug 20, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Aug 20, 2026
@coderabbitai

coderabbitai Bot commented Aug 20, 2026

Copy link
Copy Markdown
Contributor

Review Change Stack

📝 Walkthrough

Summary by CodeRabbit

  • New Features

    • Added support for creating implicit one-dimensional hierarchies with dynamic grid, cluster, and block extents.
    • Added a public hierarchy construction helper for multidimensional configurations.
    • Cooperative group types now initialize from an explicitly supplied hierarchy.
  • Improvements

    • Group construction provides clearer, more predictable hierarchy configuration.
    • One-dimensional and multidimensional launch configurations are supported and validated consistently.
  • Tests

    • Added coverage for hierarchy dimensions, return behavior, error guarantees, and supported launch configurations.

Walkthrough

Changes

The change adds implicit_hierarchy_1d() with dynamic one-dimensional extents, renames __implicit_hierarchy(), updates Cooperative Groups construction, and adds CUDA coverage for hierarchy types, extents, and launches.

Implicit hierarchy support

Layer / File(s) Summary
Hierarchy contracts and helpers
cudax/include/cuda/experimental/__group/fwd.cuh, cudax/include/cuda/experimental/__group/implicit_hierarchy.cuh
Adds __implicit_hierarchy_1d_t, renames __implicit_hierarchy() to implicit_hierarchy(), and adds implicit_hierarchy_1d() with dynamic first-axis extents.
Group construction wiring
cudax/include/cuda/experimental/__group/this_group.cuh, cudax/test/coop/any_of/*.cu, cudax/test/group/make_this_group.cu, cudax/test/group/this_group.cu
Removes implicit base construction. Cooperative Groups constructors now initialize with implicit_hierarchy(). Tests use explicit hierarchy configurations.
Hierarchy API validation
cudax/test/CMakeLists.txt, cudax/test/group/implicit_hierarchy.cu
Registers tests for return types, noexcept, runtime extents, one-dimensional and multidimensional launches, and supported cluster configurations.

Suggested reviewers: andralex

Merge Risk: ⚪ Minimal · up to 194d6

The PR exposes the hierarchy APIs and removes default construction for these groups; the remaining concerns are limited to test assertions, headers, naming, and const-correctness. No actionable merge-blocking risk remains after normal checks and review.


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

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🧹 Nitpick comments (3)
cudax/include/cuda/experimental/__group/implicit_hierarchy.cuh (2)

25-26: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Remove the unused span forward-declaration header. This file does not use span.

As per coding guidelines, “Remove unneeded headers.” As per path instructions, cudax/** changes apply CCCL common style guidance.

Sources: Coding guidelines, Path instructions


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

suggestion: Use snake_case for both local type aliases.

  • cudax/include/cuda/experimental/__group/implicit_hierarchy.cuh#L55-L55: Rename _Exts1D to a snake_case alias.
  • cudax/test/group/make_this_group.cu#L26-L26: Rename Hierarchy to hierarchy_type.

As per coding guidelines, “Use snake_case for all other symbols, except that the CUB public API uses PascalCase.”

Source: Coding guidelines

cudax/test/group/this_group.cu (1)

275-283: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick win

suggestion: Declare the test-only group values const.

  • cudax/test/group/this_group.cu#L275-L283: Change test_common_properties to accept const Group&, then declare group as const.
  • cudax/test/group/make_this_group.cu#L31-L32: Declare group as const; sync() is const and no mutation occurs.

As per coding guidelines, “All variables that are not modified must be declared const.”

Source: Coding guidelines


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 1c974cb3-f8c7-43e5-95df-88b579d1afd2

📥 Commits

Reviewing files that changed from the base of the PR and between 0619e15 and 194d633.

📒 Files selected for processing (9)
  • cudax/include/cuda/experimental/__group/fwd.cuh
  • cudax/include/cuda/experimental/__group/implicit_hierarchy.cuh
  • cudax/include/cuda/experimental/__group/this_group.cuh
  • cudax/test/CMakeLists.txt
  • cudax/test/coop/any_of/this_thread.cu
  • cudax/test/coop/any_of/this_warp.cu
  • cudax/test/group/implicit_hierarchy.cu
  • cudax/test/group/make_this_group.cu
  • cudax/test/group/this_group.cu

Included review availability: Your plan provides up to 12 included reviews per hour; 10 remain after this review.

Comment thread cudax/test/group/this_group.cu
__device__ void operator()(const Config& config)
{
test_group(cudax::this_thread{});
test_group(cudax::this_thread{cudax::implicit_hierarchy()});

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should this also test implicit_hierarchy_1d?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't think so to be honest. We will change these tests anyway sooner or later

@davebayer
davebayer enabled auto-merge (squash) August 20, 2026 08:50
@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 21m: Pass: 100%/63 | Total: 19h 59m | Max: 49m 12s | Hits: 25%/128243

See results here.

@davebayer
davebayer merged commit db98ec4 into NVIDIA:main Aug 20, 2026
93 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Aug 20, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

2 participants