Repository navigation
Conversation
MythreyaK
force-pushed
the
users/vkale/kernel-replay-rocprofv3
branch
5 times, most recently
from
September 2, 2026 05:22
65f4797 to
929884f
Compare
An error occurred while trying to automatically change base from
users/vkale/kernel-replay-rocprofv3
to
users/vkale/experimental-kernel-replay
September 2, 2026 08:40
MythreyaK
changed the base branch from
users/vkale/kernel-replay-rocprofv3
to
develop
September 2, 2026 19:27
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
This comment was marked as outdated.
This comment was marked as outdated.
MythreyaK
force-pushed
the
users/mkuriche/kernel-replay-perf-experiments
branch
from
October 1, 2026 00:44
8d63daa to
7ed8e7d
Compare
This was referenced Oct 1, 2026
Draft
1 task
This was referenced Oct 6, 2026
MythreyaK
force-pushed
the
users/mkuriche/kernel-replay-perf-experiments
branch
from
October 27, 2026 22:19
4d15c71 to
410dbb5
Compare
Contributor
|
1 task
MythreyaK
force-pushed
the
users/mkuriche/kernel-replay-perf-experiments
branch
from
November 4, 2026 21:55
41badca to
b003233
Compare
Hold the allocation inventory lock until the restore blit and dependent replay dispatch complete, preventing concurrent frees without adding host-side synchronization. Co-authored-by: Cursor Agent <cursoragent@cursor.com>
Flatten discontiguous snapshot regions into one compact 16-byte work stream so restore avoids per-tile descriptor searches and skewed grids. Add deterministic coverage for cross-region assignments and partial items. Co-authored-by: Cursor <cursoragent@cursor.com>
- compose replay snapshot storage tiers Plan whole regions against bounded GPU and pinned-host budgets before capture, using a variadic fallback chain ending in a null allocator. Materialize one arena per tier and verify mixed restore ordering, allocation rollback, and memory-pressure behavior. Improvements to come in subsequent commits. This is MVP for host-fallback. Co-authored-by: Cursor <cursoragent@cursor.com>
…king The host-fallback MVP described each storage tier twice: a planning-only allocator chain chose placement, then separate GPU and pinned-host allocation functions re-applied their own admission checks, and materialization was hard-coded to two tiers. Make each tier a single storage::arena<Policy, BackingAllocator> that owns its capacity policy and its physical backing: - planning::bounded_policy handles alignment, capacity, and offsets. - One hsa_pool_allocator template replaces the separate GPU and pinned allocation paths. Arena constructor arguments are forwarded to the backing allocator unchanged. - storage::fallback_allocator chains arenas and must end in the terminal null_allocator. Placements record the tier's storage index instead of a hard-coded storage kind. - Commit allocates only tiers that received reservations, and on failure rolls back already committed tiers in reverse order. GPU headroom and host-memory limits are now applied once, when planning budgets are computed. Allocation success remains authoritative, and a failed GPU commit still replans everything onto pinned host. Placement is unchanged: whole regions go to GPU first, then pinned host. Remove the unused global logical-offset mapper, the planning-only allocator aliases, and snapshot agent fields that nothing read. Rename the capture context's pinned_agent to cpu_agent. Move planner tests out of blit_copy.cpp into a GPU-independent kernel-replay-snapshot-storage-test covering fallback, alignment, overflow, rollback, three-tier chains, and move-only backing arguments. Co-authored-by: Cursor <cursoragent@cursor.com>
bounded_policy both sized requests and read the arena's cursor, so a new layout meant rewriting every policy. Give each tier three parts: - Admission: how much of a request the tier takes, given only what is available (whole_region; prefix_chunking splits on a granule). - Suballocator: where it goes. bump_suballocator commits one backing of exactly the used size. - Backing allocator: where memory comes from. hsa_pool_allocator now takes its storage kind and retain mode at runtime. arena<Admission, Suballocator> clamps admission to what is left, and fallback re-asks a tier until it declines, so a region can span tiers. Requests carry region_offset and priority; production still sends one request per whole region, so placement is unchanged. plan_and_commit disables a tier whose commit fails and replans, replacing the hard-coded GPU-to-pinned retry. Replace the cacheable flag with retain_mode and keep both GPU and pinned arenas across replays: one high-water arena per pool, reused when big enough, otherwise freed before the larger allocation. Retained pinned memory stays pinned until replaced or the process exits. Budgets count the retained arena as available: - GPU: min(alloc max, half of VRAM, retained + 95% of free) - Pinned: min(alloc max, retained + half of MemAvailable) BabelStream, 3 x 6 GiB behind a 221 GiB VRAM ballast (6 GiB spills to pinned): only the first replay pays the ~2 s pinned allocation, so later replays drop from ~3.3 s to ~0.63 s. All-GPU runs are unchanged. Rewrite the 11 CPU-only kernel-replay-snapshot-storage-test cases: chunking, priority, fragmented tiers, property-selected admission, replan and rollback, move-only backings, near-SIZE_MAX capacities. Co-authored-by: Cursor <cursoragent@cursor.com>
Every wave of the restore blit reads its arguments and copy descriptors. They lived in the host kernarg pool, so each restore depended on PCIe read latency: a 384 MiB restore took 240-510 us depending on how long the host had been idle. Allocate them from the GPU's pool when the CPU can map it, fence and read back the writes before the packet is published, and fall back to the host kernarg pool otherwise. The same restore now takes ~180 us, cold or warm. This impacts just the blit copy kernel, not the replayed kernel's own kernarg. Co-authored-by: Minseob Shin <minseob.shin11@gmail.com> Co-authored-by: Cursor <cursoragent@cursor.com>
The replay blit handed every thread the whole descriptor list: each thread walked every region to find its items, so a restore got slower with every region added (51 ms at 65,536 regions). Cut the regions into 16 KiB tiles on the GPU instead, and copy one tile per wave: - Three small kernels (count, scan, expand) build the tile list from the region descriptors the host already writes, one per region. - The copy kernel goes straight from its wave to a tile. Lists too short to give every wave about four tiles are split into pieces down to 4 KiB, so small restores still use the whole GPU. - The copy keeps two 16-byte loads in flight per lane and uses no scratch. The four dispatches and the existing barrier go out as one packet chain. The tile list lives in a per-agent device workspace, sized at capture and reused by restores. The kernels read only their explicit arguments, because the SDK does not fill the implicit ones, and a code object whose kernels take more is rejected. In the standalone restore benchmark on MI325X, 65,536 regions restore in 2.5 ms instead of 51 ms, host work included, and large restores copy at about 1.95 TB/s. Add kernel-replay-blit-kernels-test. It dispatches the code object as raw AQL packets, as the SDK does, and checks seven layouts byte for byte, including the gaps between regions. The host blit test now checks the tile arithmetic instead of the old per-thread rotation. Co-authored-by: Cursor <cursoragent@cursor.com>
The capture also writes a digest per 16 KiB tile, kept on the GPU with its tile list. Before pass 1, one kernel hashes the live tiles, copies back the changed ones and lists them; later passes restore only that list, and a hash of the unlisted tiles after the last pass warns if the kernel wrote others. With 2-3 passes it hashes before every pass. Full restores remain for one pass, for snapshots under 64 MiB in GPU memory and 1 MiB pinned, and for kernels that last changed most tiles or tiles off the list. To override: ROCPROFILER_INTERNAL_REPLAY_RESTORE_FORCE_MODE=auto|full|hash-every-pass|hash-once Also: full restores reuse the capture's tile list (1 dispatch, not 4), pinned extents go through the blit, restores skip the liveness walk while no allocation changes, restore counters are read from mapped device memory instead of hsa_memory_copy, and snap()'s capture loop is no longer quadratic in region count. Standalone tile benchmark (MI325X), 1% of 2 GiB changed: 0.56 ms from device memory and 0.95 ms from pinned, vs 1.14 ms and 39.8 ms for full restores. Tests: GPU kernel, policy and snapshot tests; samples force hashing on. Co-authored-by: Cursor <cursoragent@cursor.com>
…_replay_* Marks them as internal, like ROCclr's __amd_rocclr_* blit kernels. Co-authored-by: Cursor <cursoragent@cursor.com>
MythreyaK
force-pushed
the
users/mkuriche/kernel-replay-perf-experiments
branch
from
November 6, 2026 04:08
b003233 to
aeb42a1
Compare
MythreyaK
marked this pull request as ready for review
November 6, 2026 04:23
Contributor
There was a problem hiding this comment.
🟡 Changes recommended
Unresolved critical issues can leave modified inputs in later replay passes, deadlock memory-free callbacks, and allow resource reuse while GPU work is still outstanding.
7 open issues
What changed in this PR
This PR optimizes the kernel replay service introduced in #7960, reducing snapshot and restore overhead between profiling passes.
Changes:
- Adds GPU-local and pinned-host snapshot arenas with reusable storage and allocation-liveness tracking.
- Introduces batched GPU copy kernels, reusable blit resources, and hash-based restore policies.
- Expands unit tests, GPU correctness tests, bandwidth checks, and replay samples.
Paths are relative to projects/rocprofiler-sdk/.
| File | Description |
|---|---|
source/lib/rocprofiler-sdk/kernel_replay/tests/snapshot_storage.cpp |
Tests storage planning, fallback, alignment, and ownership. |
source/lib/rocprofiler-sdk/kernel_replay/tests/snap_restore.cpp |
Tests arena capture, restore, and allocation liveness. |
source/lib/rocprofiler-sdk/kernel_replay/tests/snap_bandwidth.cpp |
Updates snapshot timing and adds GPU-arena bandwidth coverage. |
source/lib/rocprofiler-sdk/kernel_replay/tests/restore_policy.cpp |
Tests restore selection and kernel history. |
source/lib/rocprofiler-sdk/kernel_replay/tests/CMakeLists.txt |
Registers new host and GPU tests. |
source/lib/rocprofiler-sdk/kernel_replay/tests/blit_kernels.cpp |
Validates GPU copies, digests, and selective restores. |
source/lib/rocprofiler-sdk/kernel_replay/tests/blit_copy.cpp |
Tests descriptor layout and tile arithmetic. |
source/lib/rocprofiler-sdk/kernel_replay/snapshot-storage.hpp |
Adds composable arenas and fallback allocation. |
source/lib/rocprofiler-sdk/kernel_replay/snapshot-plan.hpp |
Adds prioritized placement and commit retries. |
source/lib/rocprofiler-sdk/kernel_replay/restore_policy.hpp |
Defines restore modes and history interfaces. |
source/lib/rocprofiler-sdk/kernel_replay/restore_policy.cpp |
Implements restore selection and history tracking. |
source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.hpp |
Exposes generation-aware inventory updates. |
source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.cpp |
Tracks allocation-inventory changes. |
source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.hpp |
Defines arena-backed snapshots and batch-copy interfaces. |
source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.cpp |
Implements storage reuse, capture, and guarded restoration. |
source/lib/rocprofiler-sdk/kernel_replay/CMakeLists.txt |
Builds and installs blit code objects. |
source/lib/rocprofiler-sdk/kernel_replay/blit-copy.hpp |
Declares blit capture and restore operations. |
source/lib/rocprofiler-sdk/kernel_replay/blit-copy.hip |
Implements GPU tile-copy and hashing kernels. |
source/lib/rocprofiler-sdk/kernel_replay/blit-copy.cpp |
Loads kernels and manages reusable dispatch resources. |
source/lib/rocprofiler-sdk/kernel_replay/blit-copy-kernel.hpp |
Defines shared kernel arguments and tile helpers. |
source/lib/rocprofiler-sdk/hsa/queue.cpp |
Integrates batched capture and selective restores into replay. |
samples/kernel_replay/CMakeLists.txt |
Adds samples with forced hashed-restore modes. |
Review effort: Balanced
Comment on lines
+1281
to
+1285
| process_packet_batch(packets_arr, | ||
| 1, | ||
| forward_to_writer, | ||
| /*is_replay_pass=*/true, | ||
| replay_dispatch_id); |
Comment on lines
+569
to
+570
| m_impl->status = (value == 0) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR; | ||
| m_impl->completed = true; |
Comment on lines
+412
to
+413
| if(auto retained = take_retained_block(pool, size); retained) | ||
| return owned_block_t{retained, kind, pool, retain}; |
| // With two or three passes, hashing before each costs no more than reusing a list and checking | ||
| // it, and does not assume the kernel writes the same tiles every pass. | ||
| if(inputs.total_passes != 0 && inputs.total_passes <= 3) return restore_mode::hash_every_pass; | ||
| return restore_mode::hash_once; |
| PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm | ||
| PATH_SUFFIXES bin llvm/bin REQUIRED) | ||
|
|
||
| set(kernel_replay_blit_targets ${GPU_TARGETS}) |
Comment on lines
+216
to
+220
| for(const auto* directory : | ||
| {ROCPROFILER_KERNEL_REPLAY_BLIT_BUILD_DIR, ROCPROFILER_KERNEL_REPLAY_BLIT_INSTALL_DIR}) | ||
| { | ||
| auto path = fmt::format("{}/{}", directory, filename); | ||
| if(access(path.c_str(), R_OK) == 0) return path; |
A region that does not fit what is left of the GPU budget is now cut on a 16 KiB tile boundary and the rest goes to pinned host memory, instead of moving the whole region there. For example, three 6 GiB buffers under a 15 GiB GPU budget used to place 12 GiB on the GPU and 6 GiB in pinned memory. Now the first 3 GiB of the third buffer stay on the GPU, so captures and restores move half as much over PCIe. Cutting fills the GPU arena to its budget, which exposed a drift: the budget was retained + 95% of free, so holding the arena raised the next budget every replay until no VRAM was left, and each growth reallocated the arena. Budgets now take the margin from retained + free and never drop below the retained arena, which a smaller snapshot reuses whole. Tests: a cut region restores byte for byte through both restore paths, and the budget stays where it was across replays. Co-authored-by: Cursor <cursoragent@cursor.com>
Each 8-byte word of a tile is now digested with splitmix64's finalizer
keyed by its position. The old term gave some different words the same
value, for example {b-1, 0} and {0xffffffff, b} at the start of a tile,
so a tile whose only change was such a word was never restored. The
digest is still a plain sum, so tiles are still digested in pieces.
Full restores are now picked only when they move fewer bytes than
hashing: from half the tiles changed when hashing before every pass, and
from 1 - 1/(2N - 3) of them when hashing once over N passes (80% with 4
passes, 86% with 5, never without a pass count). Before, any kernel that
changed half its tiles switched to full restores for good.
A failed blit load is now cached per agent and what it created is
released, so the failure is reported and paid for once.
Co-authored-by: Cursor <cursoragent@cursor.com>
amd-vkale
requested changes
Nov 7, 2026
amd-vkale
left a comment
amd-vkale
left a comment
Contributor
There was a problem hiding this comment.
Reviewed c5724afd.
The approach is the right one. Snapshotting into a GPU arena, restoring with a blit queued ahead of the pass, and copying back only the tiles a pass changed is what makes replay cheap enough to use. The MI325X numbers in the description match that: a 4096³ DGEMM call goes from 799 ms to 11.0 ms, and the Jacobi run from 8.80 s to 0.38 s. The digest in this commit is splitmix64's finalizer keyed by word position, so a one-word change can no longer leave a tile unrestored, and 2–3 pass replays hash before every pass instead of trusting a list.
Three things still in this head can corrupt a replay or the next one. The description already lists the first two as pre-merge work, and they are not fixed yet.
I did not re-run the suite.
| if(retain == retain_mode::across_replays) | ||
| { | ||
| if(auto retained = take_retained_block(pool, size); retained) | ||
| return owned_block_t{retained, kind, pool, retain}; |
Contributor
There was a problem hiding this comment.
A retained block is returned here, before hsa_amd_agents_allow_access, and retained_backing is keyed only by pool.handle.
The GPU pool is per agent, so that is fine for the device tier. The pinned tier is the CPU pool, and GPUs on one socket share it. GPU0 allocates the arena and grants access to itself. GPU1 later takes that same block because it is large enough, and never gets access. The blit then faults on GPU1.
Grant access for the agent that is about to use the block, including this retained path, and key the retained arena by agent as well as pool so one GPU cannot be handed another's.
| HSA_WAIT_STATE_BLOCKED); | ||
| } | ||
|
|
||
| m_impl->status = (value == 0) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR; |
Contributor
There was a problem hiding this comment.
After 60 s this marks the packet completed and returns an error. ~implementation then clears resource_in_use, so the next blit on this agent reuses the kernarg buffer, the tile workspace, and this completion signal while the timed-out dispatch may still be writing them. finish() stores 1 into that same signal.
On a timeout, leave this agent's blit resources claimed and decline further replay on it. Returning an error is not enough, because the claim is dropped as soon as this packet_info is destroyed, including on the way out of a fatal restore failure.
| ROCP_INFO << fmt::format("kernel replay: {} of {} tiles changed by a pass", | ||
| counts->changed, | ||
| counts->tiles); | ||
| if(counts->unlisted_changed > 0) |
Contributor
There was a problem hiding this comment.
This runs after every pass of this dispatch has already been submitted, drained, and recorded. With the default for 4 or more passes (hash_once), a tile written off the list was not restored before the later passes, so those passes ran on modified inputs and their counter records are already delivered. The warning and the sticky history.unlisted flag only change how the next dispatch of this kernel restores.
For the dispatch that just ran, an unlisted change has to fail the replay. A warning after the records are out does not recover them. hash_every_pass for 2–3 passes does not have this hole; the longer replays, which is the rocprofv3 case this PR measures with 5 groups, do.
amd-vkale
self-requested a review
November 7, 2026 06:58
amd-vkale
approved these changes
Nov 7, 2026
amd-vkale
left a comment
amd-vkale
left a comment
Contributor
There was a problem hiding this comment.
Approving after discussion with MythreyaK understanding some optimization and implementation/tweaks still need to be done.
This branch has not been deployed
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
Kernel replay (#7960) restores device memory before every pass with blocking host copies from a host-memory snapshot, so each replayed dispatch costs hundreds of milliseconds, even for kernels that run in a millisecond or two. This PR snapshots and restores on the GPU and restores only what the kernel changed, so a replay costs little more than its passes.
Technical Details
ROCPROFILER_INTERNAL_REPLAY_RESTORE_FORCE_MODE=auto|full|hash-every-pass|hash-onceoverrides the choice.kernel-replay-snapshot-storage-test.__kernel_replay_*) are compiled withhipcc --gencoforGPU_TARGETS. Without a code object for the GPU, replay falls back to pinned-host snapshots and host copies.Assumptions:
mallocheap during replay aren't supported: they hang until replay's 60 s timeout aborts the process.Still to do before merge:
Issue Tracking
Test Plan
kernel-replay-blit-kernels-test(byte-for-byte restores over several region layouts, digests, changed-tile lists, the check for unlisted writes),kernel-replay-restore-policy-test,kernel-replay-snapshot-storage-test, and snapshot tests for module variables, cut regions and budgets.ctest -R replayon MI325X, excluding the ATT, SPM and PC-sampling tests.Test Result
Against develop with the same rocprofv3 arguments (MI325X, 5 counter groups, kernel trace):
hipblasDgemmcall, all 17 dispatches replayedUpdateKernellaunches 101–110 replayed: total run timeMicrobenchmarks (MI325X), measured as each change landed:
The blit kernels move 3.9–4.4 TB/s of DRAM traffic from device memory, within 1% of the ideal byte count; BabelStream's own kernels reach 4.3–4.7 TB/s on the same GPU. Pinned copies run at 54–55 GB/s, the practical PCIe Gen5 x16 limit.
Submission Checklist