Add support for HipGraph - #4956
Open
pfultz2 wants to merge 38 commits into
Open
Conversation
Regressions detected 🔴 |
|
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #4956 +/- ##
===========================================
+ Coverage 92.89% 93.26% +0.37%
===========================================
Files 603 623 +20
Lines 32448 33097 +649
===========================================
+ Hits 30140 30866 +726
+ Misses 2308 2231 -77 🚀 New features to boost your workflow:
|
Contributor
There was a problem hiding this comment.
Pull request overview
This PR adds HIP Graph capture/replay support to the GPU backend by introducing a new hip::graph operation plus a hipgraphify pass that partitions the root module into capturable regions and replaces them with hip::graph submodules. It also adds kernel-launch packing helpers needed to inspect/patch captured kernel node arguments, along with GPU tests covering partitioning and pointer rebinding behavior.
Changes:
- Add
gpu::hipgraphifypass to extract capturable instruction runs intohip::graphsubmodules (enabled in the default GPU backend pipeline). - Introduce
hip::graphoperation that captures a submodule into a HIP graph and replays it, with pointer-rebind logic for moved inputs. - Add kernel config pack/unpack helpers and new GPU tests for both graph partitioning and pointer rebinding.
Reviewed changes
Copilot reviewed 10 out of 10 changed files in this pull request and generated 7 comments.
Show a summary per file
| File | Description |
|---|---|
| test/gpu/pack_args.cpp | Adds tests for pack_kernel_config / unpack_kernel_config pointer-slot parsing. |
| test/gpu/hipgraphify.cpp | Adds unit tests validating hipgraphify partitioning behavior and boundaries. |
| test/gpu/hip_graph.cpp | Adds integration tests ensuring captured graphs correctly rebind moved input pointers across runs. |
| src/targets/gpu/target.cpp | Wires hipgraphify{} into the default GPU backend pass pipeline. |
| src/targets/gpu/kernel.cpp | Implements pack_kernel_config / unpack_kernel_config and uses the packed launch config helper in kernel launching. |
| src/targets/gpu/include/migraphx/gpu/kernel.hpp | Declares new kernel config helpers and exposes kernel::get_function() for correlating graph nodes to code objects. |
| src/targets/gpu/include/migraphx/gpu/hipgraphify.hpp | Adds public header for the gpu::hipgraphify pass. |
| src/targets/gpu/hipgraphify.cpp | Implements the graph partitioning/extraction pass. |
| src/targets/gpu/hip_graph.cpp | Implements the hip::graph operation: capture, instantiate, replay, and rebind (patch or re-record). |
| src/targets/gpu/CMakeLists.txt | Adds new HIP graph sources to the GPU target library build. |
Comment on lines
+115
to
+120
| check_hip(hipStreamBeginCapture(stream, hipStreamCaptureModeThreadLocal), | ||
| "hipStreamBeginCapture"); | ||
| f(); | ||
| hipGraph_t g = nullptr; | ||
| check_hip(hipStreamEndCapture(stream, &g), "hipStreamEndCapture"); | ||
| return hip_graph{share(hip_graph_ptr{g})}; |
Comment on lines
+141
to
+142
| if(not output_buffer.empty() and output_buffer.size() != outputs.size()) | ||
| return; |
Comment on lines
233
to
+235
| write_literals{.max_memory = max_memory}, | ||
| hipgraphify{}, | ||
| dead_code_elimination{}, |
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Collaborator
Author
|
@copilot resolve the merge conflicts in this pull request |
Contributor
Resolved. I merged |
pfultz2
marked this pull request as ready for review
July 10, 2026 19:56
- Bind submodule parameters by sorted name so the binding survives save/load of a compiled program (parameter order is not serialized) - Rebase the cached outputs/result onto moved parameter buffers on the patch path instead of returning views of the old buffers - Follow a code_object kernel's aliased output when collecting patch targets so kernels reading an in-place-written parameter are patched - Fall back to re-recording when movable leaf ranges overlap or a captured kernel node's argument buffer cannot be parsed back - Keep the patched kernarg buffer/config alive in graph_node_patch; HIP does not document copying extra at SetParams time - Serialize evals with a mutex since program copies share graph_state - Fall back to an uncaptured run under MIGRAPHX_TRACE_EVAL - Skip capturing a run whose output-backing allocation cannot become a submodule input (out-of-bounds alias index in release builds) - Exclude gpu::quant_gemm (same rocblas path as gpu::gemm) and hip::graph itself from capture, making the pass idempotent Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DBSRz3D6A6BaEW1P7ix9Y1
test/gpu/hip_graph.cpp: - rebind_save_load: save/load round trip of the compiled program before the rebind checks, pinning the sorted-name parameter binding - rebind_output_buffer: input and output buffers move together between runs - multi_output_tuple: two outputs captured as one graph under offload copy, returning a tuple unpacked with get_tuple_elem - rebind_aliased_inputs: one buffer bound to both inputs on the capture run, then distinct buffers, pinning the overlap fallback and the re-record path - factor the captured-graph check into captured_hip_graph() test/gpu/hipgraphify.cpp: - split_by_quant_gemm: gpu::quant_gemm is a partition boundary - idempotent: applying the pass twice equals applying it once - output_root_outside_run: a run whose output roots at an allocation reachable only through views of a pre-run value is left uncaptured All pass on gfx1201; rebind_convolution needs MIGRAPHX_DISABLE_MLIR=1 here due to a pre-existing MLIR no-solutions issue unrelated to this PR. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DBSRz3D6A6BaEW1P7ix9Y1
- rebind_sliced_input: kernels consume a slice of x, so rebinding must re-apply the captured nonzero within-leaf pointer offset - mixed_output_backing: a run mixing allocation-backed and non-allocation-backed external outputs is left uncaptured - unpack_invalid_configs: null extra, unknown tag, missing size entry, and a pointer slot past the end of the packed buffer Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01DBSRz3D6A6BaEW1P7ix9Y1
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
Technical Details
Changelog Category
Add a
CHANGELOG.mdentry for any option other thanNot Applicable