Skip to content

Add Cooperative Groups examples - #486

Open
j-stephan wants to merge 7 commits into
amd-stagingfrom
jstephan/cooperative-groups
Open

Add Cooperative Groups examples#486
j-stephan wants to merge 7 commits into
amd-stagingfrom
jstephan/cooperative-groups

Conversation

@j-stephan

@j-stephan j-stephan commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

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

@j-stephan
j-stephan requested a review from zichguan-amd July 17, 2026 12:41
@j-stephan j-stephan self-assigned this Jul 17, 2026
@j-stephan
j-stephan requested review from a team as code owners July 17, 2026 12:41

@randyh62 randyh62 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.

Looks good.

Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md Outdated
Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md Outdated
Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md Outdated
Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md Outdated
Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md Outdated
Comment thread HIP-Basic/cooperative_groups_double_buffered_tile/README.md
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
Signed-off-by: Jan Stephan <jan.stephan@amd.com>
@j-stephan
j-stephan force-pushed the jstephan/cooperative-groups branch from 43a654b to d6c2be3 Compare July 23, 2026 09:15
@j-stephan j-stephan changed the title Add Cooperative Groups double buffered tile example Add Cooperative Groups examples Jul 23, 2026
@j-stephan

Copy link
Copy Markdown
Contributor Author

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

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

Copy link
Copy Markdown
Collaborator

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.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants