Skip to content

[rocshmem] Add rocshmem_hipmodule_init API for CUDA graph compatibility - #3165

Merged
abouteiller merged 8 commits into
developfrom
users/drprajap/rshmem-hipmodule-init-2741
Feb 24, 2026
Merged

abouteiller merged 8 commits into
developfrom
users/drprajap/rshmem-hipmodule-init-2741

Conversation

@drprajap

@drprajap drprajap commented Feb 10, 2026 •

Copy link
Copy Markdown
Contributor

This commit adds a new API rocshmem_hipmodule_init() that allows
user-compiled HIP modules or device bitcode modules to be initialized with
rocSHMEM device contexts. This is essential for CUDA graph compatibility.

This API replaces manual rocshmem ctx management via rocshmem_get_device_ctx()
during JIT compilation.

Resolves: #2741

Technical Details

JIRA ID

Resolves AIROCSHMEM-145

Test Plan

Functional test is added to verify module loading after successfully setting device global symbol.

Submission Checklist

This commit adds a new API rocshmem_hipmodule_init() that allows
user-compiled HIP modules or device bitcode modules to be initialized with
rocSHMEM device contexts. This is essential for CUDA graph compatibility.

This API replaces manual rocshmem ctx management via rocshmem_get_device_ctx()
during JIT compilation.

Claude Code is used to assist in this change.

Resolves: #2741

Signed-off-by: Dimple Prajapati <dimple.prajapati@amd.com>
Implementation:
- Add HipModuleInitTester class in hipmodule_init_tester.{hpp,cpp}
- Use HIPRTC for runtime kernel compilation
- Calls rocshmem_hipmodule_init() on the compiled module to initialize device symbol.
- Launches simple kernel to verify module initialization

Signed-off-by: Dimple Prajapati <dimple.prajapati@amd.com>

Copilot AI 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.

Pull request overview

Adds a new rocSHMEM host API intended to initialize a user-loaded HIP module’s ROCSHMEM_CTX_DEFAULT symbol so dynamically-loaded device code can use rocSHMEM in CUDA/HIP graph-capture-friendly workflows.

Changes:

  • Introduces public API rocshmem_hipmodule_init(hipModule_t, hipStream_t) and implements it in rocshmem_gpu.cpp.
  • Adds a new functional test (hipmodule_init) that HIPRTC-compiles a module defining ROCSHMEM_CTX_DEFAULT and attempts to initialize it.
  • Wires the new test into the functional test harness (enum, factory, and driver mapping) and links hiprtc.

Reviewed changes

Copilot reviewed 8 out of 8 changed files in this pull request and generated 9 comments.

Show a summary per file
File Description
projects/rocshmem/include/rocshmem/rocshmem.hpp Adds the new public API declaration and documentation (removes prior API).
projects/rocshmem/src/rocshmem_gpu.cpp Implements rocshmem_hipmodule_init() using HIP module symbol lookup and async memcpy.
projects/rocshmem/tests/functional_tests/hipmodule_init_tester.hpp Declares new functional tester for HIP module init scenario.
projects/rocshmem/tests/functional_tests/hipmodule_init_tester.cpp Implements HIPRTC compile/load + calls new API, then runs a simple module kernel.
projects/rocshmem/tests/functional_tests/tester.hpp Adds new TestType enum entry for hipmodule init test.
projects/rocshmem/tests/functional_tests/tester.cpp Registers and instantiates the new tester in the factory.
projects/rocshmem/tests/functional_tests/CMakeLists.txt Adds build defines for source/binary include paths and links hiprtc.
projects/rocshmem/scripts/functional_tests/driver.sh Adds CLI name → test number mapping for hipmodule_init.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
- Address review feedback
- Fix rocshmem_hipmodule_init() API because previous implementation was using
hipMemcpyFromSymbol() with legacy/default stream. During HIP graph capture,
HIP rejects operations requiring legacy stream usage due to implicit sync rules
with capturing stream. So API needs to use hipGetSymbolAddress  +
 hipMemcpyAsync() to perform device-to-device copy of device global symbol.
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
Co-authored-by: Aurelien Bouteiller <abouteil@amd.com>
@abouteiller

Copy link
Copy Markdown
Contributor

Looks like the hiprtc is missing a dependency to compile the dynamic code in the CI @drprajap

@drprajap

Copy link
Copy Markdown
Contributor Author

Looks like the hiprtc is missing a dependency to compile the dynamic code in the CI @drprajap

yes, because local testing has rocm path set in environment, this was not exposed in my testing, I have explicitly included it now

Comment thread projects/rocshmem/src/rocshmem_gpu.cpp Resolved
@omor1

omor1 commented Feb 16, 2026

Copy link
Copy Markdown
Contributor

I think I'm somewhat confused as to why the compiled hiprtc kernel defines its own ROCSHMEM_CTX_DEFAULT symbol instead of using the one from the compiled library that it should link to. I'm not especially familiar with the hiprtc APIs and how to use them though, so I'm likely simply not understanding something about the workflow.

@drprajap

Copy link
Copy Markdown
Contributor Author

I think I'm somewhat confused as to why the compiled hiprtc kernel defines its own ROCSHMEM_CTX_DEFAULT symbol instead of using the one from the compiled library that it should link to. I'm not especially familiar with the hiprtc APIs and how to use them though, so I'm likely simply not understanding something about the workflow.

Current test is not using existing rocshmem device library kernels for hiprtc compilation because that requires pulling up lots of dependencies for kernel compilation, I tried that and with existing rocshmem device API definitions, it is not minimal effort. So I chose to orchestrate ROCSHMEM_CTX_DEFAULT device side symbol definition in a minimal device kernel module compilation. The goal of this minimal unit level functional test is to verify Device module initialization path for given compiled module using existing host allocated 'ROCSHMEM_CTX_DEFAULT' symbol and see if new rocshmem_hipmodule_init API is able to get and set the device side symbols without any API compatibility issue.

@abouteiller
abouteiller merged commit 3f15534 into develop Feb 24, 2026
10 of 11 checks passed
@abouteiller
abouteiller deleted the users/drprajap/rshmem-hipmodule-init-2741 branch February 24, 2026 18:05
radhaksri pushed a commit that referenced this pull request Feb 26, 2026
…ty (#3165)

* Add rocshmem_hipmodule_init API for CUDA graph compatibility

This commit adds a new API rocshmem_hipmodule_init() that allows
user-compiled HIP modules or device bitcode modules to be initialized with
rocSHMEM device contexts. This is essential for CUDA graph compatibility.

This API replaces manual rocshmem ctx management via rocshmem_get_device_ctx()
during JIT compilation.

Claude Code is used to assist in this change.

Resolves: #2741

Signed-off-by: Dimple Prajapati <dimple.prajapati@amd.com>

* Add functional test for rocshmem_hipmodule_init API

Implementation:
- Add HipModuleInitTester class in hipmodule_init_tester.{hpp,cpp}
- Use HIPRTC for runtime kernel compilation
- Calls rocshmem_hipmodule_init() on the compiled module to initialize device symbol.
- Launches simple kernel to verify module initialization

Signed-off-by: Dimple Prajapati <dimple.prajapati@amd.com>

* Enhance hipmodule_init test for CUDA graph capture validation

- Address review feedback
- Fix rocshmem_hipmodule_init() API because previous implementation was using
hipMemcpyFromSymbol() with legacy/default stream. During HIP graph capture,
HIP rejects operations requiring legacy stream usage due to implicit sync rules
with capturing stream. So API needs to use hipGetSymbolAddress  +
 hipMemcpyAsync() to perform device-to-device copy of device global symbol.

* return ROCSHMEM_ERROR in failure

* review feedback

* Apply suggestions from code review - Execute moduleInit test in CI

Co-authored-by: Aurelien Bouteiller <abouteil@amd.com>

* Apply suggestions from code review

missing "

* add missing rocm path for hiprtc compilation

---------

Signed-off-by: Dimple Prajapati <dimple.prajapati@amd.com>
Co-authored-by: Aurelien Bouteiller <abouteil@amd.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

[Feature]: rocshmem API for compatibility with cuda.graph capture

5 participants