Skip to content

[HIPIFY][ROCM-1254] Add cooperative_groups reduce and scans support#2448

Open
g-h-c wants to merge 2 commits into
amd-developfrom
users/g-h-c/ROCM-1254_add_cooperative_groups_reduce
Open

[HIPIFY][ROCM-1254] Add cooperative_groups reduce and scans support#2448
g-h-c wants to merge 2 commits into
amd-developfrom
users/g-h-c/ROCM-1254_add_cooperative_groups_reduce

Conversation

@g-h-c

@g-h-c g-h-c commented Mar 25, 2026 •
edited
Loading

Copy link
Copy Markdown

Motivation

cooperative_groups::reduce() was recently added in HIP.

cg::inclusive_scan and exclusive_scan are scheduled to be added in 7.14: ROCm/rocm-systems#5914

Technical Details

This PR enables automatic translation of header paths accordingly:

cooperative_groups/reduce.h to cooperative_groups/hip_reduce.h
cooperative_groups/scan.h to cooperative_groups/hip_scan.h

Test Plan

To be tested locally

Test Result

TBD

Submission Checklist

@emankov emankov changed the title ROCM-1254 - add cooperative_groups::reduce() support [HIPIFY][ROCM-1254] Add cooperative_groups::reduce() support Mar 25, 2026
@emankov emankov requested a review from ranapratap55 March 25, 2026 11:39
@emankov emankov added HIP Related to ROCm HIP Runtime or CLR sync Syncing update labels Mar 25, 2026
@emankov

emankov commented Apr 1, 2026 •
edited
Loading

Copy link
Copy Markdown
Collaborator

@mangupta,

May we ask you to provide (open) the HIPIFY ticket related to ROCM-1254 with all the upcoming HIPAMD/CLR API and data types to support in HIPIFY tools?

Thank you in advance

@emankov

emankov commented May 19, 2026

Copy link
Copy Markdown
Collaborator

Hello @g-h-c,

Are you going to continue with this draft, or are we ok to close it?

@g-h-c

g-h-c commented May 20, 2026 •
edited
Loading

Copy link
Copy Markdown
Author

Hello @g-h-c,

Are you going to continue with this draft, or are we ok to close it?

@emankov
Needs to be merged, but apart from cooperative_groups::reduce() in 7.13, inclusive/exclusive_scan() are coming in 7.14 (ROCm/rocm-systems#5914). Shall we create a hipify ticket for all of those? Do I have the permissions to create it?

Thank you

@emankov

emankov commented May 20, 2026

Copy link
Copy Markdown
Collaborator

Shall we create a hipify ticket for all of those?

Yes, please: https://github.com/ROCm/HIPIFY/issues.

Do I have the permissions to create it?

Sure, https://github.com/ROCm/HIPIFY is an open-source repo.

This was referenced Jun 15, 2026
Open
Merged
@g-h-c g-h-c marked this pull request as ready for review June 19, 2026 12:40
@g-h-c g-h-c requested review from emankov and searlmc1 as code owners June 19, 2026 12:40
@g-h-c g-h-c changed the title [HIPIFY][ROCM-1254] Add cooperative_groups::reduce() support [HIPIFY][ROCM-1254] Add cooperative_groups reduce and scans support Jun 19, 2026
ranapratap55
ranapratap55 requested changes Jun 22, 2026
View reviewed changes
Comment thread src/CUDA2HIP.cpp
{"cuda_profiler_api.h", {"hip/hip_runtime_api.h", "", CONV_INCLUDE, API_RUNTIME, 0}},
{"cooperative_groups.h", {"hip/hip_cooperative_groups.h", "", CONV_INCLUDE, API_RUNTIME, 0}},
{"cooperative_groups/reduce.h", {"hip/cooperative_groups/hip_reduce.h", "", CONV_INCLUDE, API_RUNTIME, 0}},
{"cooperative_groups/scan.h", {"hip/cooperative_groups/hip_scah", "", CONV_INCLUDE, API_RUNTIME, 0}},

@ranapratap55 ranapratap55 Jun 22, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

Suggested change
{"cooperative_groups/scan.h", {"hip/cooperative_groups/hip_scah", "", CONV_INCLUDE, API_RUNTIME, 0}},
{"cooperative_groups/scan.h", {"hip/cooperative_groups/hip_scan.h", "", CONV_INCLUDE, API_RUNTIME, 0}},

Please fix the typo.

Comment thread bin/hipify-perl
subst("channel_descriptor.h", "hip\/channel_descriptor.h", "include");
subst("cooperative_groups.h", "hip\/hip_cooperative_groups.h", "include");
subst("cooperative_groups\/reduce.h", "hip\/cooperative_groups\/hip_reduce.h", "include");
subst("cooperative_groups\/scan.h", "hip\/cooperative_groups\/hip_scan.h", "include");

@ranapratap55 ranapratap55 Jun 22, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

    subst("cooperative_groups\/scan.h", "hip\/cooperative_groups\/hip_scan.h", "include");

This looks like hand edited. After fixing the C++ typo, please regenerate this file.

@emankov emankov added feature Feature request or implementation and removed sync Syncing update labels Jun 28, 2026
emankov
emankov requested changes Jun 28, 2026
View reviewed changes

@emankov emankov left a comment

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

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

What about corresponding APIs and unit tests?

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

Reviewers

@emankov emankov emankov requested changes
@ranapratap55 ranapratap55 ranapratap55 requested changes
@searlmc1 searlmc1 Awaiting requested review from searlmc1 searlmc1 is a code owner

Requested changes must be addressed to merge this pull request.

Assignees

@mangupta mangupta

@g-h-c g-h-c

Labels

feature Feature request or implementation HIP Related to ROCm HIP Runtime or CLR

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@g-h-c @emankov @ranapratap55 @mangupta