feat(rocshmem): Implement wave-level fcollect collective operation#7900
Open
kesag wants to merge 8 commits into
Open
feat(rocshmem): Implement wave-level fcollect collective operation#7900kesag wants to merge 8 commits into
kesag wants to merge 8 commits into
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
kesag
force-pushed
the
users/kesag/rocshmem-wave-fcollect
branch
from
6ed9c51 to
0687ca9
Compare
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds wavefront-level fcollect support to rocSHMEM (typed and mem variants) across IPC and GDA backends, and introduces a new functional test to validate the wave-level collective behavior.
Changes:
- Add wave-level
rocshmem_ctx_##TYPE##_fcollect_wave()plus backend implementations (IPC, GDA) and RO-backend stubs. - Add
fcollectmem_{wg,wave}plumbing through the Context interface and backends. - Add a new functional test (
FcollectWave) and hook it into the functional test driver.
Reviewed changes
Copilot reviewed 21 out of 21 changed files in this pull request and generated 9 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/rocshmem/utils/header_files_gen/COLL.py | Extends header generation to include wave-level fcollect API declarations (and new fcollectmem declarations). |
| projects/rocshmem/tests/functional_tests/tester.hpp | Registers new FcollectWave functional test type. |
| projects/rocshmem/tests/functional_tests/tester.cpp | Wires in FcollectWaveTester creation and execution path. |
| projects/rocshmem/tests/functional_tests/tester_arguments.cpp | Enables standard argument handling defaults for the new test. |
| projects/rocshmem/tests/functional_tests/fcollect_wave_tester.hpp | Adds new tester class definition for wave-level fcollect functional testing. |
| projects/rocshmem/tests/functional_tests/fcollect_wave_tester.cpp | Implements the wave-level fcollect kernel test and host-side verification. |
| projects/rocshmem/src/rocshmem_gpu.cpp | Adds device entry points for fcollect_wave and fcollectmem_{wg,wave}, and updates fcollect WG dispatch. |
| projects/rocshmem/src/reverse_offload/context_ro_tmpl_device.hpp | Renames WG fcollect entry point and adds RO wave-level fcollect stub. |
| projects/rocshmem/src/reverse_offload/context_ro_device.hpp | Extends ROContext interface for fcollectmem and wave-level fcollect APIs. |
| projects/rocshmem/src/reverse_offload/context_ro_device.cpp | Adds RO implementations for fcollectmem_wg and a stub for fcollectmem_wave. |
| projects/rocshmem/src/ipc/context_ipc_tmpl_device.hpp | Updates IPC typed fcollect WG implementation and adds typed wave fcollect dispatch. |
| projects/rocshmem/src/ipc/context_ipc_device.hpp | Extends IPCContext interface for fcollectmem and wave-level fcollect APIs. |
| projects/rocshmem/src/ipc/context_ipc_device_coll.cpp | Implements IPC fcollectmem_{wg,wave} linear algorithms and sync. |
| projects/rocshmem/src/gda/context_gda_tmpl_device.hpp | Updates GDA typed fcollect WG implementation and adds typed wave fcollect dispatch. |
| projects/rocshmem/src/gda/context_gda_device.hpp | Extends GDAContext interface for fcollectmem and wave-level fcollect APIs. |
| projects/rocshmem/src/gda/context_gda_device_coll.cpp | Implements GDA fcollectmem_{wg,wave} linear algorithms and sync. |
| projects/rocshmem/src/context.hpp | Extends base Context API surface for fcollectmem and wave-level fcollect. |
| projects/rocshmem/src/context_tmpl_device.hpp | Adds Context::fcollect_wave() dispatch path and updates WG fcollect dispatch naming. |
| projects/rocshmem/src/context_device.cpp | Adds Context::fcollectmem_{wg,wave}() dispatch paths. |
| projects/rocshmem/scripts/functional_tests/driver.sh | Adds fcollect_wave to functional test driver mapping and execution list. |
| projects/rocshmem/include/rocshmem/rocshmem_COLL.hpp | Publishes new wave-level fcollect declarations and fcollectmem declarations in the public header. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
kesag
force-pushed
the
users/kesag/rocshmem-wave-fcollect
branch
from
fe81cc8 to
09adefe
Compare
kesag
changed the title
kesag
force-pushed
the
users/kesag/rocshmem-wave-fcollect
branch
from
8b641c0 to
cb40df2
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Motivation
Develop:
nvshmemx_##TYPE##_fcollect_warp()nvshmemx_fcollectmem_warp()to provide wave-level broadcast capabilities in rocSHMEM, matching NVSHMEM's API and semantics.
Technical Details
add wave-level fcollect implementation
rocshmem_ctx_##TYPE##_fcollect_wave()androcshmem_ctx_fcollectmem_wave()add fcollectmem_wg implementation
rocshmem_ctx_fcollectmem_wg()JIRA ID
JIRA ID: AIROCSHMEM-438
Test Plan
Test Result
Submission Checklist