Skip to content

feat(rocprofiler-sdk): kernel replay rocprofv3 integration - #10439

Merged
bgopesh merged 31 commits into
users/vkale/experimental-kernel-replayfrom
users/vkale/kernel-replay-rocprofv3
Sep 2, 2026
Merged

bgopesh merged 31 commits into
users/vkale/experimental-kernel-replayfrom
users/vkale/kernel-replay-rocprofv3

Conversation

@amd-vkale

@amd-vkale amd-vkale commented Aug 19, 2026 •

Copy link
Copy Markdown
Contributor

Motivation

Sub-PR 2/2 in the kernel replay stack. #7960 lands the ROCprofiler-SDK public API and service implementation and can merge on its own. This PR adds the rocprofv3 / rocprofiler-sdk-tool integration that makes that service reachable from the command line, plus the end-to-end test harness.

Hardware exposes a limited number of counter registers per block. ROCprofiler-SDK's existing rocprofv3 answer is application replay: partition counters into groups and re-run the entire application once per group. Cost is O(G × application wall time), and the counters come from G different executions.

Kernel replay (in #7960) re-executes each dispatch once per group within a single run, restoring device memory between passes. This PR is what lets a rocprofv3 user ask for that: --kernel-replay-beta-enabled.

Stack of PRs

PR Base Commits Scope Lands alone?
#7960 develop 4 ROCprofiler-SDK public API + implementation, docs, unit tests Yes — custom tools work; rocprofv3 unchanged
#10439 (this PR) #7960 2 rocprofv3 CLI, tool wiring, integration tests After #7960

Land order: #7960 → #10439. Review the SDK service in #7960 first. This PR assumes that API and does not re-implement snap/restore, the agent lock, or the replay loop.

Technical Details

18 files, +2561 / −3. CLI, tool, tests only.

The diagrams below are the CLI-layer before/after. The SDK replay window itself is #7960 After (1 of 2) — this PR sits on top of it.

Before — application replay (what rocprofv3 does after #7960 lands, without this PR)

This is unchanged rocprofv3 on the #7960 branch. Multiple --pmc groups still take cli_multipass and the multi-pass outer loop. Each child sees one counter group. The SDK replay service from #7960 is idle because nothing sets ROCPROF_KERNEL_REPLAY.

flowchart TB
  U["<b>rocprofv3 --pmc G0 --pmc G1 --pmc G2</b><br/>no kernel-replay flag"]
  MP["<b>cli_multipass = True</b>"]
  LOOP["<b>outer loop over groups</b>"]
  R0["run(app, pass_id=1)<br/>child process 1"]
  R1["run(app, pass_id=2)<br/>child process 2"]
  R2["run(app, pass_id=3)<br/>child process 3"]
  E0["ROCPROF_COUNTERS = G0"]
  E1["ROCPROF_COUNTERS = G1"]
  E2["ROCPROF_COUNTERS = G2"]
  SVC["<b>dispatch counting service</b><br/>one profile per dispatch"]
  WI["<b>WriteInterceptor</b><br/>normal single-pass path<br/>SDK KERNEL_REPLAY idle"]
  OUT["output pass_1/ pass_2/ pass_3/"]
  COST["<b>cost = G x T_app</b><br/>3 process launches<br/>3 distinct executions"]

  U --> MP --> LOOP
  LOOP --> R0 --> E0 --> SVC
  LOOP --> R1 --> E1 --> SVC
  LOOP --> R2 --> E2 --> SVC
  SVC --> WI --> OUT
  LOOP --> COST

  classDef cli fill:#fff8e1,stroke:#f57f17,stroke-width:2px,color:#000
  classDef sdk fill:#e0f2f1,stroke:#00695c,stroke-width:2px,color:#000
  classDef bad fill:#ffebee,stroke:#c62828,stroke-width:2px,color:#000
  class U,MP,LOOP,R0,R1,R2,E0,E1,E2,OUT cli
  class SVC,WI sdk
  class COST bad
Loading

Consequence: three runs, three sets of allocations. SQ_WAVES from group 0 does not describe the same execution as GRBM_COUNT from group 2. #7960 is on the system but unused.

After — this PR: rocprofv3 drives the #7960 replay service

The CLI merges every counter group and forces cli_multipass = False, so the child-process relaunch never runs. It sets ROCPROF_KERNEL_REPLAY. The tool registers a KERNEL_REPLAY context and answers the SDK's pass-count query with the number of counter groups. The SDK loop from #7960 then drives replay; this PR only selects group i on pass i.

flowchart TB
  U["<b>rocprofv3 --pmc G0 --pmc G1 --pmc G2<br/>--kernel-replay-beta-enabled</b>"]
  MERGE["<b>merge groups, cli_multipass = False</b><br/>bypass application replay"]
  ENV["<b>ROCPROF_KERNEL_REPLAY = 1</b>"]
  RUN["<b>single application run</b>"]
  CTX["<b>tool: KERNEL_REPLAY context</b>"]
  CFG["<b>CONFIG / PHASE_ENTER</b><br/>install pass_count_cb"]
  PC["<b>kernel_replay_pass_count_callback</b><br/>returns counter-group count for this agent"]
  SDK["<b>SDK replay loop (PR 7960)</b><br/>snapshot - pass - restore<br/>agent writer lock, one completion"]
  PASS["<b>PASS / PHASE_ENTER</b><br/>publish tl_current_replay_pass = i"]
  SEL["<b>counter_dispatch_callback</b><br/>get_replay_profile(agent, i)"]
  PACK["<b>pack tid+pass into user_data</b>"]
  REC["<b>counter_record_callback</b><br/>unpack - replay_pass on record"]
  OUT["<b>single output</b><br/>one dispatch_id per dispatch<br/>replay_pass 0..N-1"]
  WIN["<b>cost = T_app + GPU time of replayed dispatches</b><br/>all groups from the SAME execution"]

  U --> MERGE --> ENV --> RUN --> CTX --> CFG --> PC --> SDK
  SDK --> PASS --> SEL --> PACK --> REC --> OUT
  SDK -.->|"loop i = 0..N-1"| PASS
  RUN --> WIN

  classDef cli fill:#fff8e1,stroke:#f57f17,stroke-width:2px,color:#000
  classDef tool fill:#e8eaf6,stroke:#283593,stroke-width:2px,color:#000
  classDef sdk fill:#e0f2f1,stroke:#00695c,stroke-width:2px,color:#000
  classDef good fill:#e8f5e9,stroke:#2e7d32,stroke-width:2px,color:#000
  class U,MERGE,ENV,RUN cli
  class CTX,CFG,PC,SEL,PACK,REC,OUT tool
  class SDK,PASS sdk
  class WIN good
Loading

Consequence: one run, one set of allocations. replay_pass identifies which group produced each record.

Side by side

Before (application replay) After (this PR)
Process launches G 1
Cost G × T_app T_app + GPU time of replayed dispatches
Allocations Differ per group Identical across groups
SDK KERNEL_REPLAY (#7960) Idle Driving the pass loop
Output pass_1/ … pass_G/ Single output, replay_pass field
dispatch_id One per group run One per logical dispatch, shared by passes
Who repeats work rocprofv3 outer loop SDK replay loop (#7960)

Usage

rocprofv3.py

  • --kernel-replay-beta-enabled — opt-in beta flag under counter-collection options.
  • Validation — fatal error when the flag is used without --pmc.
  • Env routing — ROCPROF_KERNEL_REPLAY=1.
  • Multipass bypass — merges CLI --pmc groups and input-file pmc lines into one list-of-lists, clears input-file duplication, forces cli_multipass = False.

Changes by file

rocprofv3.py (this PR; not in #7960)

config.hpp — kernel_replay from ROCPROF_KERNEL_REPLAY

tool.cpp

Tests

Test What it proves
tests/bin/kernel-replay/ Three kernels; run_inplace_saxpy() checks buffer and __device__ global grow once per launch, never per pass
tests/rocprofv3/kernel-replay/ JSON: N passes, sanity counters constant, unique counters differ, one dispatch_id per logical dispatch
tests/kernel-replay-concurrency/ Replayed hog does not cause a concurrent victim to be replayed (30d8691184 checks hog launch errors)
tests/kernel-replay-local-context/ Per-pass local context toggles

Review Guidelines

This PR is the thin layer on top of #7960: it decides how many passes and which counter group each pass collects. It does not implement the replay loop.

Suggested review order

  1. rocprofv3.py — CLI contract and multipass bypass. A mistake here silently degrades to application replay.
  2. tool.cpp pass-count and group-selection logic.
  3. tool.cpp record attribution (user_data pack/unpack).
  4. Tests — would they fail if the above broke?

Extra scrutiny

  • The bypass is a negative condition. Correctness depends on the application-replay loop never running. Nothing positively asserts that it was skipped.
  • Two sources of "number of groups." kernel_replay_pass_count_callback and get_replay_profile both read get_agent_profiles() but are separate call sites.
  • user_data is overloaded. Under replay it carries (tid, pass); otherwise just tid.
  • Multi-agent systems. Global --pmc group count and a given agent's group count need not match.

Not in scope here (raise on #7960): replay loop, snapshot coverage, queue/agent locking, graph/multi-packet, SDMA.

Key questions for reviewers

CLI

  1. Does the bypass cover every path into use_multipass, including input-file pmc combined with CLI --pmc?
  2. Should the flag be rejected with --pid / --collection-period rather than silently taking a different path?
  3. Is a single --pmc group with the flag a silent no-op (pass count 1), or should it warn?
  4. Is the beta flag name clear enough that users will not treat it as a general performance switch?

Pass ↔ group mapping

  1. On agents with different group counts, does pass i map to group i for that agent, with no wrap-around?
  2. Returning 1 when an agent has no groups: silent skip, or user-visible?
  3. If the SDK drives more passes than advertised, is wrapping the index (non-CI) better than dropping the pass?

Record attribution

  1. Is tl_current_replay_pass always read on the thread the SDK published it on?
  2. Can a non-replayed dispatch, while a replay is active elsewhere, report a stale replay_pass of 0?
  3. Is packing (tid, pass) into the 64-bit user_data slot acceptable long-term?

Testing

  1. Does run_inplace_saxpy() fail for a partial restore, or only a total failure?
  2. Should the pytest pin replay_pass vs n before merge?
  3. Is there a test that would catch a silent fallback to application replay (presence of pass_N/)?

Integration

  1. Do downstream parsers / rocprof-compute need to understand replay_pass?
  2. Should this wait for the remove-callbacks stack, given feat(rocprofiler-sdk): kernel replay service (callback tracing) #7960 plans to drop local-context toggles after it lands?

Issue Tracking

  • Parent SDK PR: #7960
  • JIRA: AIPROFSDK-68

Test Plan

  1. Build on the feat(rocprofiler-sdk): kernel replay service (callback tracing) #7960 branch.
  2. C++ tests: rocprofiler_kernel_replay, concurrency, local-context.
  3. End-to-end:
    rocprofv3 --pmc SQ_WAVES,SQ_INSTS_VALU,GRBM_COUNT \
              --pmc SQ_WAVES,SQ_INSTS_VALU,GRBM_GUI_ACTIVE \
              --pmc SQ_WAVES,SQ_INSTS_VALU,SQ_INSTS_SALU \
              --kernel-replay-beta-enabled --output-format json -- \
              <kernel-replay test binary>
  4. Negative: flag without --pmc must fail.
  5. Regression: without the flag, application replay (pass_N/) is unchanged.

Test Result

Integration tests pass on the combined stack, on an MI-300-series GPU with a Linux OS. Restore correctness verified by run_inplace_saxpy().

Unrelated CI noise is tracked on #7960.

Submission Checklist

@github-actions github-actions Bot added documentation Improvements or additions to documentation project: rocprofiler-sdk labels Aug 19, 2026
@therock-pr-bot

therock-pr-bot Bot commented Aug 19, 2026 •

Copy link
Copy Markdown

✅ All Policy Checks Passed

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass —
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ PR Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

@amd-vkale
amd-vkale changed the base branch from develop to users/vkale/experimental-kernel-replay August 19, 2026 22:44
@amd-vkale
amd-vkale marked this pull request as ready for review August 19, 2026 22:44
@amd-vkale
amd-vkale requested review from a team as code owners August 19, 2026 22:44
@amd-vkale amd-vkale self-assigned this Aug 19, 2026
Copilot AI lite review requested due to automatic review settings August 19, 2026 22:48
@amd-vkale
amd-vkale force-pushed the users/vkale/kernel-replay-rocprofv3 branch from ec009e1 to 811afba Compare August 19, 2026 22:48

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 rocprofv3 “kernel replay” integration for multi-group counter collection in a single application run, and introduces new integration tests to validate replay correctness, per-pass context toggles, and a deterministic concurrency regression.

Changes:

  • Add --kernel-replay-beta-enabled to rocprofv3 CLI and route it to SDK tool config via ROCPROF_KERNEL_REPLAY.
  • Update rocprofiler-sdk-tool counter collection to select counter groups by SDK-provided replay pass index and emit replay_pass in JSON counter records.
  • Add new kernel-replay-focused integration tests (rocprofv3 JSON validation, deterministic concurrency repro, localized context start/stop tool + app).

Reviewed changes

Copilot reviewed 14 out of 18 changed files in this pull request and generated 2 comments.

Show a summary per file
File Description
projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/validate.py Pytest validation of replay-pass structure and counter determinism in JSON output.
projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/pytest.ini Pytest configuration for the kernel-replay rocprofv3 test directory.
projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py Pytest CLI options + fixtures to load JSON and expected replay parameters.
projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt Adds generate+validate ctest pair for rocprofv3 kernel replay counter collection.
projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt Registers the new rocprofv3/kernel-replay test subdirectory.
projects/rocprofiler-sdk/tests/kernel-replay-local-context/main.cpp HIP test app used to validate localized per-pass context toggling correctness.
projects/rocprofiler-sdk/tests/kernel-replay-local-context/CMakeLists.txt Builds local-context tool+app and defines coverage matrix for service toggles.
projects/rocprofiler-sdk/tests/kernel-replay-local-context/client.cpp LD_PRELOAD tool implementing replay targeting + local start/stop across passes.
projects/rocprofiler-sdk/tests/kernel-replay-concurrency/main.cpp Deterministic HIP workload to reproduce/guard against kernarg-pool snapshot corruption.
projects/rocprofiler-sdk/tests/kernel-replay-concurrency/CMakeLists.txt Builds and runs deterministic concurrency regression test under LD_PRELOAD tool.
projects/rocprofiler-sdk/tests/kernel-replay-concurrency/client.cpp LD_PRELOAD tool orchestrating replay window and exposing synchronization handshake.
projects/rocprofiler-sdk/tests/CMakeLists.txt Registers new kernel replay integration test suites.
projects/rocprofiler-sdk/tests/bin/kernel-replay/kernel_replay.cpp Standalone HIP app with multiple kernels + restore-correctness checks for replay.
projects/rocprofiler-sdk/tests/bin/kernel-replay/CMakeLists.txt Builds the kernel-replay test app binary.
projects/rocprofiler-sdk/tests/bin/CMakeLists.txt Adds kernel-replay test app to the tests/bin build.
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp Implements kernel-replay pass-count callback, pass-index publishing, and per-pass counter-group selection + record annotation.
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp Adds kernel_replay config flag sourced from ROCPROF_KERNEL_REPLAY.
projects/rocprofiler-sdk/source/bin/rocprofv3.py Adds CLI flag, env routing, and argument merging to keep multi-group counters in a single run under kernel replay.

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

@cursor

cursor Bot commented Aug 19, 2026

Copy link
Copy Markdown

This PR is sub-PR 2/2 in the kernel replay stack and should be reviewed after #7960.

PR Base Scope
#7960 develop ROCprofiler-SDK public API + service implementation, docs, unit tests
#10439 (this PR) #7960 rocprofv3 CLI (--kernel-replay-beta-enabled), tool wiring, integration tests

Land order: #7960 → #10439

Please review SDK internals and the callback-tracing replay service in #7960 first. This PR builds on that foundation and only adds the rocprofv3 / rocprofiler-sdk-tool integration layer plus end-to-end tests.

Reviewers: @bing-ma @SrirakshaNag

@cursor

This comment has been minimized.

@amd-vkale

Copy link
Copy Markdown
Contributor Author

@copilot Fix the code for all comments in this review thread.

When a review comment includes a suggested change, apply the suggestion exactly.

Do not make changes beyond what is described in the linked review thread.

Copilot AI commented Aug 20, 2026

Copy link
Copy Markdown
Contributor

@copilot Fix the code for all comments in this review thread.

When a review comment includes a suggested change, apply ...

Addressed in 290ea4a.

@cursor

cursor Bot commented Aug 20, 2026

Copy link
Copy Markdown

Review after #7960 — the CLI/tool slice is complete for the job it claims

This PR is the stacked rocprofv3 / tool integration. It should be reviewed after #7960. Against the updated SDK base, the diff is now what a tool PR should be:

  • rocprofv3.py, tool.cpp, config.hpp
  • tests/bin/kernel-replay/ and tests/rocprofv3/kernel-replay/
  • CLI docs: using-kernel-replay-rocprofv3, the --kernel-replay-beta-enabled sections in using-rocprofv3 / CLI options / quick guide, and the tool.cpp / rocprofv3.py source-table rows

The SDK integration tests that never call rocprofv3 dropped out of this diff; they live on #7960. Docs match the flag.

Why this is still a good beta CLI

--kernel-replay-beta-enabled does one thing and does it correctly: collect every --pmc group in a single application run, with pass count = collectable groups on that dispatch's agent. That is a coherent product for the expensive / non-deterministic application-replay case. It does not need to expose every SDK knob to be reviewable.

CLI-shaped limits — these are intentional for this slice, not missing SDK work:

Please review this after #7960. Treat both as experimental/beta. Drop Not ready to Review when opening review. Same reviewers as #7960: Bing Ma (bing-ma) and Sriraksha (SrirakshaNag).

@github-actions

Copy link
Copy Markdown
Contributor

Code Coverage Report

Code Coverage Report

Tests Only

code coverage tests.png

Samples Only

code coverage samples.png

Tests + Samples

code coverage all.png

cursoragent and others added 22 commits September 2, 2026 05:18
…r SPM

Kernel replay varies only the counter group from one pass to the next. The tool's
replay callback hands the SDK a pass count derived from the counter groups and
publishes the pass index so the counter dispatch callback can pick a group; it
never toggles a context. Every other service therefore stays enabled for the whole
replay loop and instruments all of the passes, reporting them under the single
dispatch ID replay reuses -- N thread traces, or N sets of PC samples, where the
run asked for one.

The CLI accepted these combinations silently. --att and PC sampling had no check
at all, and the SPM check only looked at --pmc, so counter groups arriving from an
input file (the shape kernel replay collapses) slipped past it. Name the
conflicting services and refuse, rather than producing output that looks right but
counts each kernel once per pass.

This is a CLI restriction only. The SDK supports per-pass service control through
the local-context API, which a custom tool can drive.

Counter-only replay is untouched, as is --att with --pmc when replay is off.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
The replay CLI logic had tests for the happy paths but not for the
shapes that actually cause trouble. Adds 22 cases over the two pure
helpers.

The one worth calling out: rocprofv3 opens the PC sampling beta on the
environment variable being *set*, not on its value, so
ROCPROFILER_PC_SAMPLING_BETA_ENABLED=0 still starts the service. If the
replay conflict check tested truthiness instead of presence, that user
would be let into a replay run that silently collects a set of PC
samples per pass. Three tests pin presence-not-truthiness.

The rest cover input-file collapse: conflicts between non-adjacent jobs,
differing kernel filters, lists that differ only in order, 0 and False
not counting as unset, every differing setting being named rather than
just the first, and input files with one job per counter group at a
realistic group count.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
The assertion Copilot corrected was vacuous in its original ordering: with the
falsey value in the first job it passed whether has_set_attr tested `is not None`
or truthiness. Note the constraint so the operands are not swapped back.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
gopesh-review-replies.md is not part of the rocprofiler-sdk project tree
and does not belong at the top level of the monorepo.

Co-authored-by: Cursor <cursoragent@cursor.com>
…est script path

TheRock installed test tree does not preserve source-relative paths to
rocprofv3.py; use $<TARGET_FILE:rocprofiler-sdk::rocprofv3> instead.

Co-authored-by: Cursor <cursoragent@cursor.com>
…re them nonzero

The cross-pass constancy check gave SQ_WAVES and SQ_INSTS_VALU the same 10%
relative band as any other counter. That band cannot help these counters and
hides the failure they exist to catch: the values are summed over every
hardware instance, so a pass that loses one of N instance records is off by
1/N -- inside the band whenever N exceeds ten.
Wave and instruction counts are fixed by launch geometry and the instruction
stream, so replay against restored inputs reproduces them bit-for-bit; measured
spread is 0% across every dispatch of a five-group run. Only the GRBM_* cycle
counters need slack, and those are unique per pass so they never reach this
check.

Also assert the common counters are nonzero in every pass. A counter stuck at
zero has min == max, so _approx_equal(0, 0) held at any tolerance and a dead
counter passed silently. The other rocprofv3 counter validators already guard
this; this one did not. Scoped to the common counters rather than
EXPECTED_COUNTERS because SQ_INSTS_LDS is legitimately zero for these kernels.

Co-authored-by: Cursor <cursoragent@cursor.com>
…ncate

The pack path tested both halves of the packed slot at runtime and reported a
violation through the logger. For the thread id that test can never fire: the
value comes from gettid(), so its width is the kernel's pid_t. Assert that width
at compile time instead, which is a guarantee rather than a per-dispatch check
and fails the build on a platform where it stops holding.

The pass index keeps its runtime check. It arrives from the SDK, so there is no
compile-time bound to assert, and truncating it would put a pass that was never
collected into the record and from there into the output.

Severity is unchanged: ROCP_CI_LOG_IF stays fatal under ROCPROFILER_CI and logs
at ERROR otherwise, rather than dropping to a warning.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…ternal env var

The tool-side kernel replay switch is an implementation detail of how rocprofv3
hands the setting to the tool library. Users configure it with
--kernel-replay-beta-enabled, so document that and drop the variable name from
the reference page and the source table.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
An input file's jobs each describe a run of their own: their own output
configuration, kernel filters and ranges. Kernel replay was folding them into a
single run, keeping only the first job's settings and rejecting input files
whose jobs disagreed about anything but pmc. Whether counter groups are
collected by replaying a dispatch or by re-running the application is a property
of a run, not a reason to redefine what the jobs are, so leave them alone and
run each one as it would run without the flag.

Counter groups given on the command line still become passes of a single run,
which is what the flag is for. Also look for pmc_groups on the input file rather
than on cmd_args, where it is never set, so a job that asks for counters that
way is no longer rejected for having none.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…elds

Give replay_user_data_t's fields member initializers so the struct cannot be
left uninitialized, and declare the replay pass thread-local with an explicit
empty initializer.

Member initializers make the struct non-trivially-default-constructible, so the
is_trivial assertion becomes is_trivially_copyable -- which is the property the
memcpy pack/unpack actually depends on. For the same reason the unpack memcpy
writes through a void*, which is what -Wclass-memaccess wants to see.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
Use replay_pass_count in the tool callback payload, drop the duplicated
rocprofv3 integration doc section left by the merge, and align the
remaining docs with the renamed SDK fields.

Co-authored-by: Cursor <cursoragent@cursor.com>
…o link

Co-authored-by: Cursor <cursoragent@cursor.com>
Metric 3 in validate.py only checked the union of non-common counters
across passes, so a single pass could collect every unique counter while
the others collected none. Require exactly one per pass and that those
per-pass counters are all distinct.

Co-authored-by: Cursor <cursoragent@cursor.com>
…ow-to

The opening paragraph said the profiler snapshots the device memory the
kernel can mutate, which implies a reachability analysis that does not
happen. Every tracked coarse-grained allocation owned by the dispatch's
agent is captured, so host RAM and copy cost track the agent's whole
footprint rather than the kernel's working set. That drives the cost of
the feature, so it belongs in the opening description and not only in
the footprint note further down.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
… counters

The CSV section claimed passes are distinguished by Counter_Name, which
only holds when each --pmc group requests distinct counters. The
kernel-replay integration test deliberately repeats sanity counters such
as SQ_WAVES in every group, and those rows share both Dispatch_Id and
Counter_Name, so CSV alone cannot say which pass produced them. State the
condition and point at JSON, where replay_pass is explicit.

Also drops the parenthetical referring to an earlier prototype: the
shipped page should describe the current design, not its history.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
…etattr

Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
… flag

Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
…-beta-enabled as acknowledgement flag

Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
The --replay-mode option pushed the kernel replay test invocations and the
CLI test arguments past the line limits the formatting CI enforces. Rewrap
them with the pinned tools: clang-format 11, cmake-format 0.6.13 and black.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
--kernel-replay-beta-enabled no longer selects kernel replay; it only
acknowledges the beta. Every replay job in the benchmark still passed the flag
on its own, so they would have run application replay instead and reported it
as replay: replay_4_groups would have measured the same thing as app_replay,
and the pass sweep would have measured nothing that varies. Ask for the
strategy explicitly.

Co-authored-by: Vivek Kale <vivek.kale@amd.com>
@MythreyaK
MythreyaK force-pushed the users/vkale/kernel-replay-rocprofv3 branch from 65f4797 to 929884f Compare September 2, 2026 05:22
@MythreyaK MythreyaK changed the title feat(rocprofiler-sdk): Sub PR of kernel replay rocprofv3 integration feat(rocprofiler-sdk): kernel replay rocprofv3 integration Sep 2, 2026
@bgopesh
bgopesh merged commit c5f49ae into develop Sep 2, 2026
58 of 70 checks passed
bgopesh added a commit that referenced this pull request Sep 2, 2026
## Motivation

<!-- Explain the purpose of this PR and the goals it aims to achieve.
-->

This PR adds an experimental **kernel replay service** to
ROCprofiler-SDK: a **public callback-tracing API and its
implementation**. It is **self-contained** — a custom ROCprofiler-SDK
tool can configure `ROCPROFILER_CALLBACK_TRACING_KERNEL_REPLAY` and
collect multiple hardware-counter groups within a **single application
run** without waiting for the `rocprofv3` CLI.

The `rocprofv3` flag, tool wiring, and end-to-end CLI tests live in
**[#10439](#10439 (stacked
on this PR). They are **not** required for this PR to merge.

Hardware exposes a limited number of counter registers per block, so a
counter set larger than that cannot be collected in a single execution
of a kernel. The established workaround is to partition the counters
into groups and re-run the entire application once per group, which
multiplies process launch, runtime initialization and all host-side work
by the number of groups, and which yields counters drawn from that many
distinct executions — so any derived metric combining counters across
groups is computed from samples that were never simultaneously true.
Kernel replay instead re-executes the individual dispatch once per group
within a single run, restoring the device memory it reads between passes
so every pass sees identical inputs.

The goals of this design are to:

1. expose replay as its own `KERNEL_REPLAY` callback-tracing domain so
it composes with counter collection, kernel tracing and thread trace
instead of displacing them;
2. ask the tool for a pass count per dispatch so that the pass count can
be the counter-group count and dispatches that need nothing opt out at
no cost;
3. support early exit and open-ended replay; and
4. give tools a loop-scoped, agent-scoped way to enable or disable a
service for a single pass, which is what makes collecting a different
counter group on each pass expressible at all.

### Stack

| PR | Base | Scope | Can land alone? |
|---|---|---|---|
| **#7960 (this PR)** | `develop` | ROCprofiler-SDK public API +
implementation, conceptual/API docs, unit tests | **Yes** — custom tools
work; `rocprofv3` is unchanged |
| [#10439](#10439) | this PR |
`rocprofv3 --kernel-replay-beta-enabled`, `rocprofiler-sdk-tool` wiring,
integration tests | After this PR |
| [#10780](#10780) | this PR |
Public ABI gate, perf-helper unit tests, benchmark configuration,
conceptual docs | Merged into this PR |
| [#10781](#10781) | #10439 |
`rocprofv3` CLI edge-case tests | Merged into #10439 |
| [#10777](#10777) | #10439 |
Earlier revision of the CLI edge-case tests | Closed; superseded by
#10781 |

**Land order: #7960 → #10439.** Review SDK internals here first. #10439
is the user-facing CLI on top of this API.

What this PR does **not** change: `rocprofv3.py`,
`rocprofiler-sdk-tool/tool.cpp`. Multiple `--pmc` groups still use
**application replay** until #10439.

**Docs note:**
[`using-kernel-replay.rst`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay.rst)
and the rocprofv3 CLI options page describe the intended
`--kernel-replay-beta-enabled` UX. That flag is **implemented in
#10439**. If this PR merges first, the SDK API is live; the CLI flag
will 404 until #10439 follows. Custom tools are the supported path in
the gap.

### Application replay today (and why kernel replay exists)

**Application replay is already the default in `rocprofv3`.** When a
tool requests more counters than fit in one hardware pass, the counters
are split into groups and the **whole application is re-run from start
to finish once per group**. With multiple `--pmc` flags, `rocprofv3`
relaunches the application in a child process for each group and writes
output under `pass_N/` directories (see [`rocprofv3.py` multi-pass
loop](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/bin/rocprofv3.py#L2569)
— this file is **unchanged** by this PR).

That model is simple and correct, but its cost is **`O(G × application
wall time)`**. Kernel replay is the alternative when you need all groups
on the **same** dispatches from **one** run. This PR provides that
alternative **as an SDK service**. Opting `rocprofv3` into it is
[#10439](#10439).

### Lineage and related work

| Item | Role |
|---|---|
| Branch
[`users/amd-vkale/kernel-replay-basic-buildable`](https://github.com/ROCm/rocm-systems/tree/users/amd-vkale/kernel-replay-basic-buildable)
| **Earlier buildable prototype** where kernel replay was tied to
counter collection
([`rocprofiler_configure_kernel_replay_counting_service()`](https://github.com/ROCm/rocm-systems/blob/users/amd-vkale/kernel-replay-basic-buildable/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/kernel_replay.cpp#L34-L60),
[`--kernel-replay`](https://github.com/ROCm/rocm-systems/blob/users/amd-vkale/kernel-replay-basic-buildable/projects/rocprofiler-sdk/source/bin/rocprofv3.py#L700-L711)).
This PR replaces that interface while keeping the same core mechanism
(memory tracker, snapshot/restore, replay loop in `WriteInterceptor`). |
| [#8586](#8586) | Vision to
remove the per-queue callback registry. |
| [#8891](#8891),
[#8790](#8790),
[#8895](#8895),
[#8887](#8887) |
Remove-callbacks stack (preferred land-before order). |
| [#10193](#10193) | Product
vision: kernel replay on the callback-removal stack. **DO NOT MERGE**. |
| [#10439](#10439) |
`rocprofv3` / tool integration on top of **this** PR. |

### Land assumption (important for reviewers)

**This PR is written against current `develop` and assumes the
remove-callbacks stack does *not* land before it.** Per-pass service
selection uses **localized context toggles**
([`replay_local_start_context_cb` /
`replay_local_stop_context_cb`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h#L99-L100),
backed by
[`local_context.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/local_context.cpp#L73-L97)).

**Ideally [#8891](#8891) →
[#8790](#8790) →
[#8895](#8895) →
[#8887](#8887) land first.** If
they do, this branch rebases and the toggle layer is removed in favor of
explicit dispatch hooks — replay window, snapshot/restore, and the
callback-tracing API remain.

Suggested land order:

1. [#8891](#8891) counters
2. [#8790](#8790) thread trace
3. [#8895](#8895) PC sampling
4. [#8887](#8887) SPM
5. **This PR**
6. [#10439](#10439) `rocprofv3`
integration

## Technical Details

<!-- Explain the changes along with any relevant GitHub links. -->

**55 files, +5760 / −45.** Public API, SDK service, docs, unit tests. No
`rocprofv3` / `rocprofiler-sdk-tool` sources.

### The problem

A GPU has a fixed number of hardware counter registers per block. Every
profiler using ROCprofiler-SDK currently partitions oversized counter
sets and re-runs the **entire application** once per group. That is `G`
process launches and `G` distinct executions. Derived metrics that
combine counters across groups are computed from samples that were never
simultaneously true.

Kernel replay moves the repetition inside: re-execute the individual
*dispatch* G times within one run, restoring device memory between
passes.

### Design goals

1. **Replay composes with other services.** Replay is a property of how
a dispatch is executed, not a kind of counter collection.
2. **Pass count is per dispatch.** `pass_count_cb` returning `<= 1` opts
a dispatch out.
3. **Per-pass service selection (interim).** Localized context overrides
until the remove-callbacks stack lands.

### What replay has to guarantee

1. **Identical inputs per pass** — snap before the loop, restore between
passes.
2. **Exclusive device access for the window** — per-agent writer lock;
no other dispatch on that agent mutates captured state.
3. **Records complete before revert** — drain the async handler before
restore.
4. **One completion, one dispatch** — the application sees a single
kernel completion.
5. **Per-pass service selection** — enable a different counter group
each pass without global start/stop.

### Approach against the goals

| Goal | Mechanism | Code |
|---|---|---|
| Composes rather than displaces |
`ROCPROFILER_CALLBACK_TRACING_KERNEL_REPLAY` domain (`CONFIG` + `PASS`);
configuring it enables the memory tracker and forces AQL queue
interception | [`callback_tracing.cpp`
L168–186](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp#L168-L186),
[`memory_tracker.cpp`
L160](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.cpp#L160)
|
| Pass count per dispatch | `pass_count_cb` asked once at `CONFIG`
enter; `total_passes > 1` (or 0 + `replay_continue_cb`) requests replay
| [`replay_callbacks.cpp`
L194–221](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/replay_callbacks.cpp#L194-L221)
|
| Uninteresting dispatches cost nothing | `pass_count_cb` returning 1 →
`replay_requested = false`, fall through to the normal path |
[`replay_callbacks.cpp`
L218–221](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/replay_callbacks.cpp#L218-L221)
|
| Early exit / open-ended replay | `replay_continue_cb` at each `PASS`
exit | [`kernel_replay.h`
L90–93](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h#L90-L93)
|
| Per-pass service control (interim) |
`replay_local_start/stop_context_cb`, thread-local override map armed
for the loop | [`local_context.cpp`
L73–97](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/local_context.cpp#L73-L97)
|
| Identical inputs | `memory_snapshot::snap(agent)` / `restore()` —
coarse VRAM + module `__device__` / `__constant__` |
[`memory_snapshot.cpp`
L136–256](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.cpp#L136-L256)
|
| Exclusive access | Per-agent reader/writer lock; replay takes the
**writer**; barrier + sibling drain | [`queue.cpp`
L111–126](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L111-L126),
[`queue.cpp`
L948–990](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L948-L990)
|
| Records complete before revert | Drain async completion handlers after
each pass | [`queue.cpp`
L147–155](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L147-L155)
|
| One completion | App signal suppressed per pass; fired once after the
loop | [`queue.cpp`
L951–956](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L951-L956)
|

### Architecture

Four diagrams. The two **Before** pictures are prior art. The two
**After** pictures are this stack: **After (1 of 2) is what this PR
lands**; **After (2 of 2) is what #10439 adds on top**.

#### Before (1 of 2) — application replay (still the `rocprofv3`
default; unchanged by this PR)

`rocprofv3` detects multi-`--pmc`
([`cli_multipass`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/bin/rocprofv3.py#L2426-L2429))
and **re-executes the application** once per group ([multi-pass
loop](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/bin/rocprofv3.py#L2569)).
Each child configures ordinary dispatch counting. The SDK never replays.

```mermaid
flowchart TB
  U1["<b>rocprofv3 --pmc G0 --pmc G1 --pmc G2</b>"]
  MP1["<b>cli_multipass = True</b>"]
  LOOP1["<b>outer loop: one child process per group</b>"]
  R0["run(app, pass_id=1)"]
  R1["run(app, pass_id=2)"]
  R2["run(app, pass_id=3)"]
  SVC1["<b>dispatch counting service</b><br/>one profile per dispatch"]
  WI1["<b>WriteInterceptor</b><br/>normal single-pass path<br/>no snapshot / no restore"]
  OUT1["output pass_1/ pass_2/ pass_3/"]
  COST1["<b>cost = G x T_app</b><br/>G distinct executions"]

  U1 --> MP1 --> LOOP1
  LOOP1 --> R0 --> SVC1
  LOOP1 --> R1 --> SVC1
  LOOP1 --> R2 --> SVC1
  SVC1 --> WI1 --> OUT1
  LOOP1 --> COST1

  classDef cli fill:#fff8e1,stroke:#f57f17,stroke-width:2px,color:#000
  classDef sdk fill:#e0f2f1,stroke:#00695c,stroke-width:2px,color:#000
  classDef cost fill:#ffebee,stroke:#c62828,stroke-width:2px,color:#000
  class U1,MP1,LOOP1,R0,R1,R2,OUT1 cli
  class SVC1,WI1 sdk
  class COST1 cost
```

*This PR does not change that path.* `rocprofv3` users still get
application replay until
[#10439](#10439).

#### Before (2 of 2) — original kernel-replay counting-service prototype

The first experiment kept **one application run** but exposed replay as
[`rocprofiler_configure_kernel_replay_counting_service()`](https://github.com/ROCm/rocm-systems/blob/users/amd-vkale/kernel-replay-basic-buildable/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/kernel_replay.cpp#L34-L60)
— mutually exclusive with dispatch counter collection, global pass
count, no agent writer lock.

```mermaid
flowchart TB
  T0["<b>Tool</b><br/>rocprofv3 --kernel-replay --kernel-replay-passes N"]
  CFG["<b>configure_kernel_replay_counting_service()</b><br/>mutually exclusive with<br/>dispatch counter collection"]
  WI0["WriteInterceptor replay loop"]
  SNAP0["Snapshot::snap / restore"]
  CSV0["CSV Replay_Pass column"]

  T0 --> CFG --> WI0 --> SNAP0 --> CSV0

  classDef old fill:#fff3e0,stroke:#e65100,stroke-width:2px,color:#000
  class T0,CFG,WI0,SNAP0,CSV0 old
```

#### After (1 of 2) — this PR: SDK callback-tracing domain (no
`rocprofv3`)

This is what **lands if #7960 merges by itself.** Any ROCprofiler-SDK
client configures `KERNEL_REPLAY`, answers `pass_count_cb`, and
(optionally) toggles services per pass. The SDK owns snap/restore, the
agent lock, and the replay window. `rocprofv3` is not in this picture.

```mermaid
flowchart TB
  T2["<b>Any ROCprofiler-SDK tool</b><br/>custom client — not rocprofv3"]
  DOM["<b>KERNEL_REPLAY callback-tracing domain</b><br/>CONFIG and PASS<br/>composes with every other service"]
  ASK["<b>pass_count_cb(dispatch_info)</b><br/>per-dispatch pass count<br/>return 1 to opt this dispatch out"]
  CONT["<b>replay_continue_cb</b><br/>early exit / indefinite loop"]
  WI2["WriteInterceptor replay window"]
  LOCK["<b>per-agent writer lock</b>"]
  DRAIN["drain this queue, then siblings"]
  SNAP2["<b>snap(agent)</b><br/>coarse VRAM + module variables"]
  PASS["<b>pass loop</b><br/>PASS_ENTER - submit - drain handler - PASS_EXIT"]
  REST2["restore between passes"]
  OVR["<b>local_context override</b> (interim)"]
  SIG["application completion signal fires once"]

  T2 --> DOM --> ASK --> WI2
  DOM --> CONT --> PASS
  WI2 --> LOCK --> DRAIN --> SNAP2 --> PASS
  PASS --> REST2 --> PASS
  PASS --> OVR
  PASS --> SIG

  classDef api fill:#e3f2fd,stroke:#1565c0,stroke-width:2px,color:#000
  classDef new fill:#e8f5e9,stroke:#2e7d32,stroke-width:2px,color:#000
  classDef win fill:#f3e5f5,stroke:#6a1b9a,stroke-width:2px,color:#000
  class DOM,ASK,CONT api
  class SNAP2,OVR,PASS new
  class LOCK,DRAIN,REST2,WI2,SIG,T2 win
```

Code for this diagram (all in **this** PR):

- API:
[`experimental/kernel_replay.h`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h#L82-L150)
- Configure path: [`callback_tracing.cpp`
L168–186](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp#L168-L186)
- Pass-count delivery: [`replay_callbacks.cpp`
L194–221](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/replay_callbacks.cpp#L194-L221)
- Replay window: [`queue.cpp`
L917–1070](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L917-L1070)
- Snap/restore:
[`memory_snapshot.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.cpp#L136-L256)
- Queue intercept: [`queue_controller.cpp`
L673–712](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp#L673-L712)
- `replay_pass` on JSON records: [`counter_info.hpp`
L104–116](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/output/counter_info.hpp#L104-L116)

#### After (2 of 2) — once
[#10439](#10439) lands:
`rocprofv3` on top of this API

**Not in this PR.** #10439 merges all `--pmc` groups, sets
`ROCPROF_KERNEL_REPLAY`, and the default tool answers `pass_count_cb`
with the group count and selects group *i* on pass *i*. Review that
wiring there.

```mermaid
flowchart TB
  U["<b>rocprofv3 --pmc G0 --pmc G1 --pmc G2<br/>--kernel-replay-beta-enabled</b>"]
  MERGE["<b>merge groups, cli_multipass = False</b><br/>PR 10439 — not this PR"]
  ENV["ROCPROF_KERNEL_REPLAY = 1"]
  RUN["single application run"]
  SDK["<b>SDK replay loop (this PR)</b><br/>snapshot - pass - restore"]
  SEL["<b>tool selects group i on pass i</b><br/>PR 10439"]
  OUT["single output, replay_pass 0..N-1"]

  U --> MERGE --> ENV --> RUN --> SDK --> SEL --> OUT

  classDef cli fill:#fff8e1,stroke:#f57f17,stroke-width:2px,color:#000
  classDef sdk fill:#e0f2f1,stroke:#00695c,stroke-width:2px,color:#000
  classDef later fill:#e8eaf6,stroke:#283593,stroke-width:2px,color:#000
  class U,MERGE,ENV,RUN cli
  class SDK sdk
  class SEL,OUT later
```

Details and review questions for that diagram:
[#10439](#10439).

### API surface

New domain `ROCPROFILER_CALLBACK_TRACING_KERNEL_REPLAY`, operations
`CONFIG` and `PASS`.
[`rocprofiler_callback_tracing_kernel_replay_data_t`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h#L82-L150)
carries `pass_count_cb`, `replay_continue_cb`, `current_pass` /
`total_passes`, and the interim local-context callbacks.

Configuring the service enables the memory tracker and forces AQL queue
interception ([`queue_controller.cpp`
L673–712](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp#L673-L712)).
A second `KERNEL_REPLAY` subscriber is rejected ([`callback_tracing.cpp`
L168–170](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp#L168-L170)).

### Replay window

For a single-dispatch write with replay requested ([`queue.cpp`
L917–1070](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L917-L1070)):

1. `CONFIG PHASE_ENTER` → `pass_count_cb` → N
2. Per-agent writer lock
3. Barrier on this queue + drain siblings
4. `memory_snapshot::snap(agent)` — decline replay if incomplete
5. Pass loop: `PASS ENTER` → submit → drain async handler → `PASS EXIT`
→ restore
6. `CONFIG PHASE_EXIT`; fire the application's completion signal once

HIP graph launches are not replayed ([`queue.cpp`
L918–923](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L918-L923)):
warn once, run once.

### Solution (algorithm)

```
WriteInterceptor(queue, one dispatch packet):
    plan <- CONFIG_PHASE_ENTER(dispatch_info)
    N    <- plan.pass_count_cb(dispatch_info)
    if N <= 1 and not plan.indefinite: CONFIG_PHASE_EXIT(); return NORMAL_PATH
    if graph_launch_active: warn once; return NORMAL_PATH

    acquire_writer(agent_lock); barrier(queue); sync(sibling queues)
    snapshot <- snap(agent)
    if snapshot incomplete: CONFIG_PHASE_EXIT(); return NORMAL_PATH
    for pass in 0 .. N-1:
        armed: PASS_PHASE_ENTER(pass)
        submit(process_packet_batch(shared dispatch_id, app signal suppressed))
        wait(pass_done); drain_async_handlers(queue)
        PASS_PHASE_EXIT(pass)
        if not should_continue_replay(pass): break
        restore(snapshot)
    CONFIG_PHASE_EXIT(); submit(barrier carrying app completion signal)
    release_writer(agent_lock)
```

### Carveouts and known limitations

- **Experimental.** Header under
`include/rocprofiler-sdk/experimental/`.
- Localized context toggles are **interim** until remove-callbacks
lands.
- `hipMallocAsync` / `hipMallocManaged` carved out; `hsa_amd_vmem_map`
is not hooked.
- Host memory is out of scope. Cache state is not restored.
- `hsa_amd_memory_async_copy` is not serialized against a replay window
(SDMA TODO).
- Only single-packet, single-dispatch submissions are replayed;
multi-packet batches warn and run once.
- Single process; no MPI coordination.
- Stuck drains abort after bounded waits (~60 s) rather than hang.

## Review guidelines

This PR is the **SDK service**. Do not review `rocprofv3` CLI or
`tool.cpp` here — those files are not in the diff. Review them on
[#10439](#10439).

**Suggested review order (source and docs, not tests)**

1. **Replay loop and concurrency** — [`hsa/queue.cpp`
L917–1088](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L917-L1088),
lock at
[L111–126](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp#L111-L126),
conceptual
[`kernel_replay_concurrency_and_isolation.md`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/kernel_replay_concurrency_and_isolation.md)
2. **Callback API** —
[`experimental/kernel_replay.h`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h),
[`replay_callbacks.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/replay_callbacks.cpp),
single-subscriber rule in [`callback_tracing.cpp`
L168–186](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp#L168-L186)
3. **Memory snapshot** —
[`memory_snapshot.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.cpp),
[`memory_tracker.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_tracker.cpp)
4. **Local context toggles** —
[`local_context.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/local_context.cpp),
consulted from counters / SPM / thread trace dispatch handlers

### Key questions for reviewers

1. Is the writer-lock scope (whole window, not per pass) the right
isolation vs throughput trade-off?
2. Are drain-then-fatal paths acceptable for a beta feature?
3. Is graph / multi-packet graceful decline (warn + run once) the right
contract?
4. Is omitting executable-flag pools from the snapshot acceptable for
beta?
5. Is declining replay on a partial snapshot (host OOM) the right
fallback vs aborting the process?
6. Are the local-context override semantics (what each service honors)
documented clearly enough for a tool author using this PR **without**
`rocprofv3`?
7. Should the how-to docs that mention `--kernel-replay-beta-enabled`
move to #10439 so this PR's docs match what actually ships if it merges
first?

## Issue Tracking

JIRA ID: AIPROFSDK-68

Follow-on CLI: [#10439](#10439)

## Test Plan

<!-- Explain any relevant testing done to verify this PR. -->

### Unit tests (this PR)

| CTest target | Source | Coverage |
|---|---|---|
| `kernel-replay-snapshot-test` |
[`snap_restore.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_restore.cpp)
| Capture, restore, in-place accumulation, module variables, pool filter
|
| `kernel-replay-local-context-test` |
[`local_context.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/local_context.cpp)
| Arm window, sticky overrides, loop teardown |
| `kernel-replay-configure-test` |
[`replay_configure.cpp`](https://github.com/ROCm/rocm-systems/blob/users/vkale/experimental-kernel-replay/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/replay_configure.cpp)
| Single-subscriber rule |

Service-level local-context tests: `counters/tests/local_context.cpp`,
`spm/tests/local_context.cpp`, `thread_trace/tests/local_context.cpp`,
`pc_sampling/tests/local_context.cpp`.

Filter: `ctest -R kernel-replay` (unit tests in this PR).

### Integration / rocprofv3 tests — **#10439, not this PR**

`tests/bin/kernel-replay/`, `tests/rocprofv3/kernel-replay/`,
`tests/kernel-replay-concurrency/`, `tests/kernel-replay-local-context/`
moved with the CLI.

## Test Result

<!-- Briefly summarize test outcomes. -->

Kernel-replay **unit** tests pass on the ROCprofiler-SDK GPU matrix
(Core • mi325, gfx94X). End-to-end `rocprofv3
--kernel-replay-beta-enabled` results belong to
[#10439](#10439). Unrelated CI
noise (hip-stream-reuse ordering, coverage SEGFAULT, MIOpen/RCCL matrix)
is tracked on this PR's discussion and is not kernel-replay-specific.

## Submission Checklist

- [ ] Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
- [ ] This PR is reviewable without #10439.
- [ ] Remove `Not ready to Review` before requesting final review.

---------

Co-authored-by: Mythreya <mythreya.kuricheti@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Gopesh Bhardwaj <gopesh.bhardwaj@amd.com>
Co-authored-by: Jonathan R. Madsen <jrmadsen@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

amd-vkale pushed a commit to amd-vkale/rocm-systems that referenced this pull request Oct 9, 2026
… toggles

Kernel replay (ROCm#7960) shipped an interim per-pass service mask for the
per-queue callback registry era: replay_start_context / replay_stop_context
on the PASS payload, backed by a thread-local override map
(kernel_replay/local_context) that counters, SPM, thread trace and kernel
dispatch tracing each consulted in their dispatch paths. ROCm#7960 planned to
drop it once the remove-callbacks stack landed and every service ran from an
explicit queue hook.

With the hooks in place, each pass is a separate submit on the replaying
thread, and every dispatch-scoped service already asks the tool's own
dispatch callback what to collect on it: counting and SPM through the
counter configuration, dispatch thread trace through its control flags.
The mask duplicated that decision, could not cover the agent-wide services
(a PC sampling or device counting toggle reported success and did nothing),
and needed a matching check in every consumer. Remove it:

- kernel_replay.h: drop replay_start_context, replay_stop_context and the
  rocprofiler_kernel_replay_context_cb_t typedef; document selecting
  services per pass from their dispatch callbacks.
- Delete kernel_replay/local_context.{hpp,cpp} and stop arming it in
  replay_callbacks.cpp and installing it around the pass loop in
  hsa/queue.cpp.
- counters, SPM and thread trace dispatch handlers go back to their own
  enabled checks; kernel dispatch tracing reports every pass.
- Delete the override unit tests (kernel_replay, counters, SPM, thread
  trace, PC sampling). The unrelated serialization-owner test that shared
  thread_trace/tests/local_context.cpp moves to serialization_test.cpp
  unchanged.
- ABI test: pin the shorter record, which now ends at replay_continue.
- Samples: att, spm, service-sequence and counters-then-pc-sampling select
  their per-pass services in the dispatch callbacks from the pass
  published at PASS PHASE_ENTER, like the counters sample already did.
- Docs and CHANGELOG updated.

rocprofv3 (ROCm#10439) never used the toggles: it already selects counter
group i on pass i in its dispatch counting callback, so its behavior is
unchanged.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

documentation Improvements or additions to documentation organization: ROCm project: rocprofiler-sdk

Projects

None yet

Development

Successfully merging this pull request may close these issues.

8 participants