Repository navigation
fix(rocprofiler-sdk): decline range replay when kernel arguments point outside the snapshot - #14030
Open
minseobshin11 wants to merge 1 commit into
Conversation
Open
2 tasks
minseobshin11
marked this pull request as ready for review
October 22, 2026 18:44
2 tasks
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
Range replay restores the coarse-grained device snapshot before every pass. The documented scope already excludes unified/managed, host, fine-grained and virtual-memory-mapped allocations, and said such ranges are "not detected — see the divergence check". But the divergence check hashes only snapshot regions, so it cannot see this memory. Those ranges replayed with
REPLAYED, and a kernel writing that memory applied its writes once per pass. A counter incremented 4 times in the application's run read 10 after 3 passes, with divergence 0.Technical Details
When a dispatch is recorded inside a range, its explicit kernel arguments are checked. The trailing 256-byte HIP hidden-argument block is skipped. An argument that points at GPU-writable memory outside the snapshot declines the range with the existing
SNAPSHOT_FAILED, and a warning names the memory kind. This covers host-pinned, managed, locked, fine-grained, another agent's memory, virtual-memory mappings, and pageable host memory when the runtime reports XNACK enabled. The application's own run is unchanged.What the opt-in divergence check (
ROCPROF_RANGE_REPLAY_VERIFY=1) catches on the PR head, compared with this guard:hipHostMallocmemoryhipMallocAsyncmemoryhipMemcpyAsyncinto pinned host memoryhipMallocbufferThe first six rows were measured with the check on. The others follow from what each check reads. The guard acts before any replayed pass. The divergence check reports after the passes, when the application's result is already wrong, so it remains the diagnostic for what the guard cannot see.
No new public status, environment variable, kernel replay change, or memory tracker change. The
SNAPSHOT_FAILEDcomment, the header warning and the conceptual docs describe the check.Device-side
mallocreplays correctly: HIP's device heap is allocated through the tracked allocator, so the snapshot covers it. Growing the heap inside a range declinesALLOCATION_CHANGED_IN_RANGEthrough the existing allocation check.Test Plan
range-replay-untracked-memory-test: memory classification, hidden-argument skip, and decline decisions.hipMalloc,hipHostMalloc,hipMallocManaged,hipMallocAsync, fine-grained and uncached device memory, another GPU's fine-grained memory, smallhipMemcpyAsyncin both directions, devicemallocwithin and beyond the initial heap, and OpenMP zero-copy under XNACK.Test Result
hipMallocmallocwithin HIP's heapmallocthat grows the heapALLOCATION_CHANGED_IN_RANGE, correcthipHostMalloc,hipMallocManaged,hipMallocAsynchipMemcpyAsync(4 KiB, both directions, pageable and pinned)mallocctest -j8on the PR head too and pass when rerun alone.dot, which now declines because its partial-sum buffer ishipHostMallocmemory. Every application's own check passes.Submission Checklist