Add Cooperative Groups examples by j-stephan · Pull Request #486 · ROCm/rocm-examples

j-stephan · 2026-07-17T12:41:02Z

Motivation

In ROCm 7.14, we introduced new Cooperative Groups features. This example demonstrates how to use memcpy_async and the new inclusive / exclusive scan functionality.

Technical Details

Not supported for HIP SDK or CUDA.

Test Plan

Tested locally on gfx1100 lab machine.

Test Result

Works.

Added/Updated documentation?

Yes
No, does not apply to this PR.

Included Visual Studio files?

Yes
No, does not apply to this PR.

Submission Checklist

New examples contain proper CMake and Make files.
I have updated relevant CI workflows and the new examples are being built and tested in CI.
Check and guard against unsupported ASICs if relevant dependent libraries have limited support.
Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

randyh62

Looks good.

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

j-stephan · 2026-07-23T09:16:47Z

@randyh62 @sourabhuday - I've added a new prefix_sum example to this PR. Please review those files as well.

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

…-examples into jstephan/cooperative-groups

zichguan-amd · 2026-07-27T14:02:44Z

hip/cooperative_groups/hip_scan.h doesn't seem to be in the 7.14 apt packages. https://rocm.docs.amd.com/projects/HIP/en/docs-7.14.0/how-to/hip_runtime_api/cooperative_groups.html#unsupported-nvidia-cuda-features, looks like develop has it but not released in 7.14.

zichguan-amd · 2026-07-30T14:21:56Z

Just need to use SKIP_FROM_BUILD similar to https://github.com/ROCm/rocm-examples/blob/d6017851e7d0cb860b6a78b8d2016735a31dcc01/Libraries/hipFFT/Makefile to skip the build. I haven't found a better solution that can inject this info without modifying all the Makefiles.

zichguan-amd

LGTM if CI passes

randyh62 · 2026-08-03T18:21:44Z

+A single block streams a 1D array through two LDS (shared memory) buffers, one tile at a time.
+The async load of the next tile is issued into the second buffer while the current tile is consumed.
+A split barrier separates the moment a thread has finished reading a buffer from the moment the
+block guarantees that every thread is done. The kernel applies the element-wise operation
+`out[i] = scale * in[i] + bias`, which is trivial to validate against a CPU reference.


Is there a reference to the docs that can be added here?

randyh62 · 2026-08-03T23:34:57Z

+run as independent work between `barrier_arrive` and `barrier_wait`. Correctness (no data races,
+validated output) is the top priority.
+
+This example targets the AMD/HIP (ROCm) backend, and it requires a ROCm version recent enough to ship `hip/cooperative_groups/memcpy_async.h`.


You could mention a specific version, or later than...

jujiang-del · 2026-08-05T21:05:31Z

+This program showcases a double-buffered tile load pipeline built from two cooperative groups
+APIs: the group-collective `cooperative_groups::memcpy_async` and the split barrier of a
+`thread_block`. A split barrier decomposes an ordinary block barrier into two phases,
+`barrier_arrive` and `barrier_wait`, so that independent work can run between them instead of every
+thread blocking immediately.


Suggested change

This program showcases a double-buffered tile load pipeline built from two cooperative groups

APIs: the group-collective `cooperative_groups::memcpy_async` and the split barrier of a

`thread_block`. A split barrier decomposes an ordinary block barrier into two phases,

`barrier_arrive` and `barrier_wait`, so that independent work can run between them instead of every

thread blocking immediately.

This program demonstrates double-buffered tile load pipeline built with two cooperative groups

APIs:

* `cooperative_groups::memcpy_async`

* The `thread_block` split barrier (`barrier_arrive()` / `barrier_wait()`)

A split barrier decomposes an ordinary block barrier into two phases:

* `barrier_arrive`

* `barrier_wait`

This separation allows useful independent work to execute between the two operations instead of forcing every thread to block immediately at a single barrier.

jujiang-del · 2026-08-05T21:21:03Z

+`cooperative_groups::memcpy_async` is an **asynchronous**, group-collective copy (typically
+global <-> LDS). HIP does not expose a separate wait handle (there is no `cg::wait()`), so its
+completion must be enforced by a following group barrier - either a `block.sync()` (as the official
+reference test does) or, as in this example, the `barrier_wait` of a split barrier whose
+`barrier_arrive` is issued *after* the copy. Ordering matters: the prefetch of the next tile is


Suggested change

`cooperative_groups::memcpy_async` is an **asynchronous**, group-collective copy (typically

global <-> LDS). HIP does not expose a separate wait handle (there is no `cg::wait()`), so its

completion must be enforced by a following group barrier - either a `block.sync()` (as the official

reference test does) or, as in this example, the `barrier_wait` of a split barrier whose

`barrier_arrive` is issued *after* the copy. Ordering matters: the prefetch of the next tile is

`cooperative_groups::memcpy_async()` is an **asynchronous**, group-collective copy operation, typically used to stage data between global memory and LDS (shared memory). HIP does not provide a separate completion handle or equivalent cg::wait() primitive, Therefore, completion of an asynchronous copy must be synchronized through a subsequent group barrier. The official reference test uses block.sync(), while this example relies on the `barrier_wait()` phase of a split barrier.

Correct ordering is critical. The prefetch of the next tile is

jujiang-del

See my comments.

j-stephan requested a review from zichguan-amd July 17, 2026 12:41

j-stephan self-assigned this Jul 17, 2026

j-stephan requested review from a team as code owners July 17, 2026 12:41

randyh62 approved these changes Jul 17, 2026

View reviewed changes

j-stephan mentioned this pull request Jul 20, 2026

docs(hip): Add link to Cooperative Groups double-buffered tile example ROCm/rocm-systems#8846

Open

1 task

sourabhuday self-requested a review July 20, 2026 21:00

sourabhuday approved these changes Jul 20, 2026

View reviewed changes

j-stephan added 4 commits July 23, 2026 10:19

Add Cooperative Groups double buffered tile example

bf7ae47

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

Incorporate feedback

cdc5fd5

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

CMake cleanup

707ee73

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

Add Cooperative Groups prefix sum example

d6c2be3

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

j-stephan force-pushed the jstephan/cooperative-groups branch from 43a654b to d6c2be3 Compare July 23, 2026 09:15

j-stephan changed the title ~~Add Cooperative Groups double buffered tile example~~ Add Cooperative Groups examples Jul 23, 2026

j-stephan requested review from randyh62 and sourabhuday July 23, 2026 09:16

zichguan-amd reviewed Jul 24, 2026

View reviewed changes

Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/CMakeLists.txt Outdated

Comment thread HIP-Basic/cooperative_groups_prefix_sum/CMakeLists.txt Outdated

zichguan-amd and others added 3 commits July 25, 2026 23:13

Merge branch 'amd-staging' into jstephan/cooperative-groups

65f9ed7

Incorporate feedback

7ad9c8a

Signed-off-by: Jan Stephan <jan.stephan@amd.com>

Merge branch 'jstephan/cooperative-groups' of github-public:ROCm/rocm…

17f774e

…-examples into jstephan/cooperative-groups

Merge branch 'amd-staging' into jstephan/cooperative-groups

9860413

zichguan-amd approved these changes Jul 31, 2026

View reviewed changes

sourabhuday approved these changes Jul 31, 2026

View reviewed changes

randyh62 approved these changes Aug 4, 2026

View reviewed changes

jujiang-del reviewed Aug 5, 2026

View reviewed changes

-This program showcases a double-buffered tile load pipeline built from two cooperative groups
-APIs: the group-collective `cooperative_groups::memcpy_async` and the split barrier of a
-`thread_block`. A split barrier decomposes an ordinary block barrier into two phases,
-`barrier_arrive` and `barrier_wait`, so that independent work can run between them instead of every
-thread blocking immediately.
+This program demonstrates double-buffered tile load pipeline built with two cooperative groups
+APIs:
+* `cooperative_groups::memcpy_async`
+* The `thread_block` split barrier (`barrier_arrive()` / `barrier_wait()`)
+A split barrier decomposes an ordinary block barrier into two phases:
+* `barrier_arrive`
+* `barrier_wait`
+This separation allows useful independent work to execute between the two operations instead of forcing every thread to block immediately at a single barrier.

Uh oh!

Conversation

j-stephan commented Jul 17, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Test Plan

Test Result

Added/Updated documentation?

Included Visual Studio files?

Submission Checklist

Uh oh!

randyh62 left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

j-stephan commented Jul 23, 2026

Uh oh!

Uh oh!

Uh oh!

zichguan-amd commented Jul 27, 2026

Uh oh!

zichguan-amd commented Jul 30, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

zichguan-amd left a comment

Choose a reason for hiding this comment

Uh oh!

randyh62 Aug 3, 2026

Choose a reason for hiding this comment

Uh oh!

randyh62 Aug 3, 2026

Choose a reason for hiding this comment

Uh oh!

jujiang-del Aug 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jujiang-del Aug 5, 2026

Choose a reason for hiding this comment

Uh oh!

jujiang-del left a comment

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

j-stephan commented Jul 17, 2026 •

edited

Loading

zichguan-amd commented Jul 30, 2026 •

edited

Loading

jujiang-del Aug 5, 2026 •

edited

Loading