test(hrr): [AIRUNTIME-2665] reach the remaining _spt APIs via the compile option by andraghetti · Pull Request #10306 · ROCm/rocm-systems

andraghetti · 2026-08-18T11:09:15Z

Motivation

_spt is not an API surface applications call by name. It is what the ordinary
entry points resolve to when a translation unit is built with
-fgpu-default-stream=per-thread. The HRR suite still reached 28 of those entry
points by writing hipMemset_spt(...) and friends explicitly, which no
application does and which tests a calling convention the runtime does not see in
the field.

#8840 moved the four per-thread memsets to a translation unit built with that
option and had them call the ordinary names. This applies the same treatment to
the remaining _spt calls, so hrr_workload_test.cc no longer names an _spt
entry point anywhere.

Technical Details

Test-only change. No runtime or generator code is touched.

The workloads move into hrr_spt_workload_test.cc, the unit #8840 added, and
call the ordinary names. Two conventions carry the design:

Oracle policy. A replay-validated buffer has to be written and read by APIs
that both snapshot a blob at capture and reproduce their effect at replay. Where
the API under test is a no-op on replay it operates on a scratch buffer and the
validated buffer is written separately, so a deliberate no-op cannot be mistaken
for a regression. Which of the two applies is not a judgement call, it follows
NOOP_PLAYBACK_APIS:

Workload	APIs under test	Coverage
`MemcpySpt`	`hipMemcpy`, `hipMemcpyAsync`, `hipMemcpy2D`, `hipMemcpy2DAsync`, `hipMemcpyToSymbol`, `hipMemcpyToSymbolAsync`	capture path (all six are `NOOP_PLAYBACK`)
`Memset3DSpt`	`hipMemset3D`, `hipMemset3DAsync`	capture path (`NOOP_PLAYBACK`, as are the plain 3-D memsets)
`Memcpy3DSpt`	`hipMemcpy3D`, `hipMemcpy3DAsync`	replay effect, one validated destination per API
`StreamQuerySpt`	`hipStreamIsCapturing`, `hipStreamQuery`, `hipStreamSynchronize`, `hipStreamGetPriority`, `hipStreamGetFlags`, `hipEventRecord`	recorded ids plus a replay that completes; these write no device memory
`StreamWaitEventSpt`	`hipStreamWaitEvent`	replay effect, through cross-stream ordering
`StreamCaptureBeginSpt`	`hipStreamBeginCapture`	replay effect, through a graph that executes
`GraphLaunchSpt`	`hipGraphLaunch`	replay effect, through a graph that executes
`StreamCaptureQuerySpt`	`hipStreamIsCapturing`, `hipStreamGetCaptureInfo`	replay effect, through a graph that executes

Since the source no longer names the API under test, the archive is the only
evidence the redirect took effect, so every roundtrip driver asserts the recorded
ids before replaying. Losing the per-source option would otherwise leave a
capture that records the plain entry points and replays perfectly well, passing
while testing something already covered elsewhere.

Macro ordering. The workloads disagree about which names must stay
redirected: one tests hipMemcpy_spt, the rest need the plain hipMemcpy as
their oracle, and the graph frame needs the plain stream-capture pair. HRR
replays a graph only when it can find it in graph_map, and only the
hand-written hipStreamEndCapture handler puts it there, so a frame closed
through the redirected name leaves hipGraphInstantiate with nothing to
instantiate. The redirects are therefore undone in stages, hipStreamEndCapture
before hipStreamBeginCapture so the workload whose API under test is the
begin call still reaches it. Each stage is followed by a preprocessor check
rather than a comment asking the next reader to be careful, so reordering the
file breaks the build instead of quietly testing the plain entry point. The
StreamCaptureBeginSpt roundtrip asserts the _spt begin id and the plain end
id together, which is what proves the staging did what it claims.

The patterns are repeated bytes (0x41, 0x42, 0x44, 0x45, 0x46, 0x4A,
0x4C, 0x5A) for the reason #8840 documented: an expected value has to decode
above atol/(1-rtol) in every encoding the D2H validator tries before an
all-zero replay buffer is reported as a mismatch. Unit_HRR_StreamCaptureBeginSpt
previously used 0x33, which decodes to 4.17e-8 as f32 and so passed against
zeros under the default tolerance; the roundtrips now also pin
HIP_HRR_D2H_EXACT=1.

hrr_run_exact_roundtrip() takes an optional list of expected API ids, and
hrr_require_recorded_apis() generalises the id assertions #8840 added to
Unit_HRR_MemsetSptRoundtrip.

Issue Tracking

JIRA ID: AIRUNTIME-2665

Test Plan

Build libamdhip64, hrr-playback and the HrrTest Catch2 target on a gfx950
host, then run the full HRR unit suite as a regression pass.

Test Result

Verified on a gfx950 host at this revision: the full HRR unit suite passes, both
the GPU-free pass (19/19) and the GPU pass (52/52), with
Unit_HRR_MemcpyPeerRoundtrip skipped as it needs a second GPU. All nine
per-thread roundtrips pass, the four added here and the five that already
existed. The four workloads whose _spt blocks were removed
(MemsetExtra, MemcpyExtra, StreamAdvanced2, HostRegLaunch) still pass.

The build is itself part of the evidence for the macro staging: the guard before
the memcpy workload fails the build if hipMemcpy is no longer redirected there,
and the guard after the #undef fails it if the redirect survived, so a clean
compile pins the macro state at each stage.

Submission Checklist

Look over the contributing guidelines at
https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

therock-pr-bot · 2026-08-18T11:10:18Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	PR does not contain code files — Unit Test auto-passed
🚫 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?

Copilot

Pull request overview

This PR refactors HRR “per-thread default stream” coverage by moving remaining explicit _spt API call sites out of the shared hrr_workload_test.cc into a dedicated translation unit compiled with per-thread-default-stream enabled, and registers the new direct workloads in the HRR unit config.

Changes:

Removed explicit _spt call sites from hrr_workload_test.cc (and updated related workload comments accordingly).
Added hrr_spt_workload_test.cc containing new hidden [.][hrr-direct] workloads intended to exercise _spt entry points via compile-option macro redirection.
Updated HRR unit build/config to compile the new TU with per-thread default stream and to register the new direct workload names in hrr.yaml.

Reviewed changes

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

File	Description
projects/hip-tests/catch/unit/hrr/hrr_workload_test.cc	Removes explicit `_spt` workloads/coverage blocks and updates workload descriptions.
projects/hip-tests/catch/unit/hrr/hrr_spt_workload_test.cc	Adds new per-thread-default-stream direct workloads compiled under per-thread default stream.
projects/hip-tests/catch/unit/hrr/CMakeLists.txt	Adds the new TU and applies per-source per-thread-default-stream compile options.
projects/hip-tests/catch/config/configs/unit/hrr.yaml	Registers the new direct workload names for HRR unit configuration.

Suppressed comments (2)

projects/hip-tests/catch/unit/hrr/hrr_spt_workload_test.cc:262

This workload is intended to run the *_spt capture-query APIs inside a replay-supported stream-capture frame. With the per-thread compile option, hipStreamBeginCapture/hipStreamEndCapture are macro-redirected to the *_spt variants; playback_hipStreamEndCapture_spt does not populate graph_map (projects/clr/hipamd/src/hrr/playback/hip_playback_generated.cpp:3735-3740), so hipGraphInstantiate will fail on replay. Temporarily undef the begin/end macros so the capture frame uses the non-spt APIs, while leaving the query calls redirected to *_spt.

  HIP_CHECK(hipStreamBeginCapture(s, hipStreamCaptureModeThreadLocal));

  hipStreamCaptureStatus st = hipStreamCaptureStatusNone;
  HIP_CHECK(hipStreamIsCapturing(s, &st));
  REQUIRE(st == hipStreamCaptureStatusActive);

projects/hip-tests/catch/unit/hrr/hrr_spt_workload_test.cc:309

Under the per-thread compile option, hipStreamEndCapture is macro-redirected to hipStreamEndCapture_spt. HRR playback's *_spt EndCapture handler does not record the returned graph into graph_map (projects/clr/hipamd/src/hrr/playback/hip_playback_generated.cpp:3735-3740), so the subsequent hipGraphInstantiate will fail on replay. Keep BeginCapture redirected to *_spt (as intended for this workload) but call the non-spt hipStreamEndCapture so replay can translate the graph handle (projects/clr/hipamd/src/hrr/playback/hip_playback.cpp:2072-2095).

  hipGraph_t g = nullptr;
  HIP_CHECK(hipStreamEndCapture(s, &g));
  hipGraphExec_t exec = nullptr;

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

The graph workloads keep the plain stream-capture pair: only the hand-written hipStreamEndCapture handler records a graph that replay can instantiate, so a frame closed through the redirected name leaves hipGraphInstantiate with nothing to find.

…hrr-spt-test-conversion

Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>

#11157 has landed, so the HRR standalone-project layout is now in develop and this branch's tests move with it into projects/hrr/tests/integration. Resolution notes: * hrr_workload_test.cc — this branch's removal of the _spt exercise blocks is kept in all seven conflicting regions. They were converted into standalone hidden [.][hrr-direct] cases in hrr_spt_workload_test.cc, which is where develop's copies of them would otherwise be duplicated. No _spt call remains in hrr_workload_test.cc. * hrr_roundtrip_test.cc — this branch's four new _spt roundtrip cases are kept. * The relocated sources are rewritten onto the HRR harness: HIP_TEST_CASE -> HRR_TEST_CASE (5) and HIP_CHECK -> HRR_HIP_CHECK (122 across the workload and SPT translation units). * The 32 hrr.yaml entries and the catch/unit/hrr/CMakeLists.txt edit are dropped in favour of projects/hrr/tests/integration/CMakeLists.txt, which already builds hrr_spt_workload_test.cc with the per-thread-default-stream compile option this conversion depends on. Develop's hip-tests capture sentinel is taken unchanged.

#10221 has merged, so this PR's base auto-retargeted from its branch to develop. Because #10221 was squash-merged, its commits on this branch have no ancestry in common with develop's copy, which rendered its seven files as changes here. Merging develop settles that: the diff is back to the generator and its three generated outputs. #10306 has also merged in the meantime and arrives with the same merge.

…pile option (#10306) ## Motivation `_spt` is not an API surface applications call by name. It is what the ordinary entry points resolve to when a translation unit is built with `-fgpu-default-stream=per-thread`. The HRR suite still reached 28 of those entry points by writing `hipMemset_spt(...)` and friends explicitly, which no application does and which tests a calling convention the runtime does not see in the field. #8840 moved the four per-thread memsets to a translation unit built with that option and had them call the ordinary names. This applies the same treatment to the remaining `_spt` calls, so `hrr_workload_test.cc` no longer names an `_spt` entry point anywhere. ## Technical Details Test-only change. No runtime or generator code is touched. The workloads move into `hrr_spt_workload_test.cc`, the unit #8840 added, and call the ordinary names. Two conventions carry the design: **Oracle policy.** A replay-validated buffer has to be written and read by APIs that both snapshot a blob at capture and reproduce their effect at replay. Where the API under test is a no-op on replay it operates on a scratch buffer and the validated buffer is written separately, so a deliberate no-op cannot be mistaken for a regression. Which of the two applies is not a judgement call, it follows `NOOP_PLAYBACK_APIS`: | Workload | APIs under test | Coverage | |---|---|---| | `MemcpySpt` | `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpy2D`, `hipMemcpy2DAsync`, `hipMemcpyToSymbol`, `hipMemcpyToSymbolAsync` | capture path (all six are `NOOP_PLAYBACK`) | | `Memset3DSpt` | `hipMemset3D`, `hipMemset3DAsync` | capture path (`NOOP_PLAYBACK`, as are the plain 3-D memsets) | | `Memcpy3DSpt` | `hipMemcpy3D`, `hipMemcpy3DAsync` | replay effect, one validated destination per API | | `StreamQuerySpt` | `hipStreamIsCapturing`, `hipStreamQuery`, `hipStreamSynchronize`, `hipStreamGetPriority`, `hipStreamGetFlags`, `hipEventRecord` | recorded ids plus a replay that completes; these write no device memory | | `StreamWaitEventSpt` | `hipStreamWaitEvent` | replay effect, through cross-stream ordering | | `StreamCaptureBeginSpt` | `hipStreamBeginCapture` | replay effect, through a graph that executes | | `GraphLaunchSpt` | `hipGraphLaunch` | replay effect, through a graph that executes | | `StreamCaptureQuerySpt` | `hipStreamIsCapturing`, `hipStreamGetCaptureInfo` | replay effect, through a graph that executes | Since the source no longer names the API under test, the archive is the only evidence the redirect took effect, so every roundtrip driver asserts the recorded ids before replaying. Losing the per-source option would otherwise leave a capture that records the plain entry points and replays perfectly well, passing while testing something already covered elsewhere. **Macro ordering.** The workloads disagree about which names must stay redirected: one tests `hipMemcpy_spt`, the rest need the plain `hipMemcpy` as their oracle, and the graph frame needs the plain stream-capture pair. HRR replays a graph only when it can find it in `graph_map`, and only the hand-written `hipStreamEndCapture` handler puts it there, so a frame closed through the redirected name leaves `hipGraphInstantiate` with nothing to instantiate. The redirects are therefore undone in stages, `hipStreamEndCapture` before `hipStreamBeginCapture` so the workload whose API under test *is* the begin call still reaches it. Each stage is followed by a preprocessor check rather than a comment asking the next reader to be careful, so reordering the file breaks the build instead of quietly testing the plain entry point. The `StreamCaptureBeginSpt` roundtrip asserts the `_spt` begin id and the plain end id together, which is what proves the staging did what it claims. The patterns are repeated bytes (`0x41`, `0x42`, `0x44`, `0x45`, `0x46`, `0x4A`, `0x4C`, `0x5A`) for the reason #8840 documented: an expected value has to decode above `atol/(1-rtol)` in every encoding the D2H validator tries before an all-zero replay buffer is reported as a mismatch. `Unit_HRR_StreamCaptureBeginSpt` previously used `0x33`, which decodes to 4.17e-8 as f32 and so passed against zeros under the default tolerance; the roundtrips now also pin `HIP_HRR_D2H_EXACT=1`. `hrr_run_exact_roundtrip()` takes an optional list of expected API ids, and `hrr_require_recorded_apis()` generalises the id assertions #8840 added to `Unit_HRR_MemsetSptRoundtrip`. ## Issue Tracking JIRA ID: AIRUNTIME-2665 ## Test Plan Build `libamdhip64`, `hrr-playback` and the `HrrTest` Catch2 target on a gfx950 host, then run the full HRR unit suite as a regression pass. ## Test Result Verified on a gfx950 host at this revision: the full HRR unit suite passes, both the GPU-free pass (19/19) and the GPU pass (52/52), with `Unit_HRR_MemcpyPeerRoundtrip` skipped as it needs a second GPU. All nine per-thread roundtrips pass, the four added here and the five that already existed. The four workloads whose `_spt` blocks were removed (`MemsetExtra`, `MemcpyExtra`, `StreamAdvanced2`, `HostRegLaunch`) still pass. The build is itself part of the evidence for the macro staging: the guard before the memcpy workload fails the build if `hipMemcpy` is no longer redirected there, and the guard after the `#undef` fails it if the redirect survived, so a clean compile pins the macro state at each stage. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <113136742+danielbolsson@users.noreply.github.com>

… into users/danielbolsson/hrr-regions Carries the parent branch down the stack. It brings develop with it, now containing the merged #10221 and #10306, plus the review fix on #10632 that corrects the reserved-slot comment in gen_hrr_api_args.py. No conflicts. Re-running the generator over the merged tree produces an empty diff, so the checked-in artifacts remain generator output.

… into users/danielbolsson/hrr-regions-memcpy Carries the parent branch down the stack, bringing develop (now with #10221 and #10306 merged) and the #10632 review fix with it. No conflicts; the generator reproduces the checked-in artifacts unchanged.

…memcpy' into users/danielbolsson/hrr-test-harness Carries the parent branch down the stack, bringing develop with it. One conflict, in hrr_roundtrip_test.cc, and it is not the simple one it looks like. This branch deletes the file's local hrr_capture_direct and hrr_playback_env because they move into hrr_test_common.hh, so the obvious resolution is to keep the deletion wholesale. That would have been wrong: #10306 has since merged into develop and added hrr_require_recorded_apis to the same block. It has no counterpart in the shared header and is still called by Unit_HRR_*SptRoundtrip below, so deleting the block would have removed a definition that is still used. Resolved by keeping only hrr_require_recorded_apis from the incoming side and dropping the two helpers the shared header now supplies. Every helper the file references resolves to either the header or a surviving local definition, and none is defined twice.

…ness' into users/danielbolsson/hrr-api-matrix Carries the parent branch down the stack, bringing develop with it. No conflicts: the API-matrix tests, the watchdog rebuilt on SpawnProc::runWithTimeout, and hrr_require_recorded_apis (kept from #10306 when resolving the parent) all merge unchanged. The generator reproduces the checked-in artifacts.

systems-assistant · 2026-09-10T09:31:24Z

TheRock Submodule Bump Activity

Newest first

2026-09-10 — Included in TheRock via ROCm/TheRock#8099.

…pile option (ROCm#10306) ## Motivation `_spt` is not an API surface applications call by name. It is what the ordinary entry points resolve to when a translation unit is built with `-fgpu-default-stream=per-thread`. The HRR suite still reached 28 of those entry points by writing `hipMemset_spt(...)` and friends explicitly, which no application does and which tests a calling convention the runtime does not see in the field. ROCm#8840 moved the four per-thread memsets to a translation unit built with that option and had them call the ordinary names. This applies the same treatment to the remaining `_spt` calls, so `hrr_workload_test.cc` no longer names an `_spt` entry point anywhere. ## Technical Details Test-only change. No runtime or generator code is touched. The workloads move into `hrr_spt_workload_test.cc`, the unit ROCm#8840 added, and call the ordinary names. Two conventions carry the design: **Oracle policy.** A replay-validated buffer has to be written and read by APIs that both snapshot a blob at capture and reproduce their effect at replay. Where the API under test is a no-op on replay it operates on a scratch buffer and the validated buffer is written separately, so a deliberate no-op cannot be mistaken for a regression. Which of the two applies is not a judgement call, it follows `NOOP_PLAYBACK_APIS`: | Workload | APIs under test | Coverage | |---|---|---| | `MemcpySpt` | `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpy2D`, `hipMemcpy2DAsync`, `hipMemcpyToSymbol`, `hipMemcpyToSymbolAsync` | capture path (all six are `NOOP_PLAYBACK`) | | `Memset3DSpt` | `hipMemset3D`, `hipMemset3DAsync` | capture path (`NOOP_PLAYBACK`, as are the plain 3-D memsets) | | `Memcpy3DSpt` | `hipMemcpy3D`, `hipMemcpy3DAsync` | replay effect, one validated destination per API | | `StreamQuerySpt` | `hipStreamIsCapturing`, `hipStreamQuery`, `hipStreamSynchronize`, `hipStreamGetPriority`, `hipStreamGetFlags`, `hipEventRecord` | recorded ids plus a replay that completes; these write no device memory | | `StreamWaitEventSpt` | `hipStreamWaitEvent` | replay effect, through cross-stream ordering | | `StreamCaptureBeginSpt` | `hipStreamBeginCapture` | replay effect, through a graph that executes | | `GraphLaunchSpt` | `hipGraphLaunch` | replay effect, through a graph that executes | | `StreamCaptureQuerySpt` | `hipStreamIsCapturing`, `hipStreamGetCaptureInfo` | replay effect, through a graph that executes | Since the source no longer names the API under test, the archive is the only evidence the redirect took effect, so every roundtrip driver asserts the recorded ids before replaying. Losing the per-source option would otherwise leave a capture that records the plain entry points and replays perfectly well, passing while testing something already covered elsewhere. **Macro ordering.** The workloads disagree about which names must stay redirected: one tests `hipMemcpy_spt`, the rest need the plain `hipMemcpy` as their oracle, and the graph frame needs the plain stream-capture pair. HRR replays a graph only when it can find it in `graph_map`, and only the hand-written `hipStreamEndCapture` handler puts it there, so a frame closed through the redirected name leaves `hipGraphInstantiate` with nothing to instantiate. The redirects are therefore undone in stages, `hipStreamEndCapture` before `hipStreamBeginCapture` so the workload whose API under test *is* the begin call still reaches it. Each stage is followed by a preprocessor check rather than a comment asking the next reader to be careful, so reordering the file breaks the build instead of quietly testing the plain entry point. The `StreamCaptureBeginSpt` roundtrip asserts the `_spt` begin id and the plain end id together, which is what proves the staging did what it claims. The patterns are repeated bytes (`0x41`, `0x42`, `0x44`, `0x45`, `0x46`, `0x4A`, `0x4C`, `0x5A`) for the reason ROCm#8840 documented: an expected value has to decode above `atol/(1-rtol)` in every encoding the D2H validator tries before an all-zero replay buffer is reported as a mismatch. `Unit_HRR_StreamCaptureBeginSpt` previously used `0x33`, which decodes to 4.17e-8 as f32 and so passed against zeros under the default tolerance; the roundtrips now also pin `HIP_HRR_D2H_EXACT=1`. `hrr_run_exact_roundtrip()` takes an optional list of expected API ids, and `hrr_require_recorded_apis()` generalises the id assertions ROCm#8840 added to `Unit_HRR_MemsetSptRoundtrip`. ## Issue Tracking JIRA ID: AIRUNTIME-2665 ## Test Plan Build `libamdhip64`, `hrr-playback` and the `HrrTest` Catch2 target on a gfx950 host, then run the full HRR unit suite as a regression pass. ## Test Result Verified on a gfx950 host at this revision: the full HRR unit suite passes, both the GPU-free pass (19/19) and the GPU pass (52/52), with `Unit_HRR_MemcpyPeerRoundtrip` skipped as it needs a second GPU. All nine per-thread roundtrips pass, the four added here and the five that already existed. The four workloads whose `_spt` blocks were removed (`MemsetExtra`, `MemcpyExtra`, `StreamAdvanced2`, `HostRegLaunch`) still pass. The build is itself part of the evidence for the macro staging: the guard before the memcpy workload fails the build if `hipMemcpy` is no longer redirected there, and the guard after the `#undef` fails it if the redirect survived, so a clean compile pins the macro state at each stage. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <113136742+danielbolsson@users.noreply.github.com>

…pile option (#10306) ## Motivation `_spt` is not an API surface applications call by name. It is what the ordinary entry points resolve to when a translation unit is built with `-fgpu-default-stream=per-thread`. The HRR suite still reached 28 of those entry points by writing `hipMemset_spt(...)` and friends explicitly, which no application does and which tests a calling convention the runtime does not see in the field. #8840 moved the four per-thread memsets to a translation unit built with that option and had them call the ordinary names. This applies the same treatment to the remaining `_spt` calls, so `hrr_workload_test.cc` no longer names an `_spt` entry point anywhere. ## Technical Details Test-only change. No runtime or generator code is touched. The workloads move into `hrr_spt_workload_test.cc`, the unit #8840 added, and call the ordinary names. Two conventions carry the design: **Oracle policy.** A replay-validated buffer has to be written and read by APIs that both snapshot a blob at capture and reproduce their effect at replay. Where the API under test is a no-op on replay it operates on a scratch buffer and the validated buffer is written separately, so a deliberate no-op cannot be mistaken for a regression. Which of the two applies is not a judgement call, it follows `NOOP_PLAYBACK_APIS`: | Workload | APIs under test | Coverage | |---|---|---| | `MemcpySpt` | `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpy2D`, `hipMemcpy2DAsync`, `hipMemcpyToSymbol`, `hipMemcpyToSymbolAsync` | capture path (all six are `NOOP_PLAYBACK`) | | `Memset3DSpt` | `hipMemset3D`, `hipMemset3DAsync` | capture path (`NOOP_PLAYBACK`, as are the plain 3-D memsets) | | `Memcpy3DSpt` | `hipMemcpy3D`, `hipMemcpy3DAsync` | replay effect, one validated destination per API | | `StreamQuerySpt` | `hipStreamIsCapturing`, `hipStreamQuery`, `hipStreamSynchronize`, `hipStreamGetPriority`, `hipStreamGetFlags`, `hipEventRecord` | recorded ids plus a replay that completes; these write no device memory | | `StreamWaitEventSpt` | `hipStreamWaitEvent` | replay effect, through cross-stream ordering | | `StreamCaptureBeginSpt` | `hipStreamBeginCapture` | replay effect, through a graph that executes | | `GraphLaunchSpt` | `hipGraphLaunch` | replay effect, through a graph that executes | | `StreamCaptureQuerySpt` | `hipStreamIsCapturing`, `hipStreamGetCaptureInfo` | replay effect, through a graph that executes | Since the source no longer names the API under test, the archive is the only evidence the redirect took effect, so every roundtrip driver asserts the recorded ids before replaying. Losing the per-source option would otherwise leave a capture that records the plain entry points and replays perfectly well, passing while testing something already covered elsewhere. **Macro ordering.** The workloads disagree about which names must stay redirected: one tests `hipMemcpy_spt`, the rest need the plain `hipMemcpy` as their oracle, and the graph frame needs the plain stream-capture pair. HRR replays a graph only when it can find it in `graph_map`, and only the hand-written `hipStreamEndCapture` handler puts it there, so a frame closed through the redirected name leaves `hipGraphInstantiate` with nothing to instantiate. The redirects are therefore undone in stages, `hipStreamEndCapture` before `hipStreamBeginCapture` so the workload whose API under test *is* the begin call still reaches it. Each stage is followed by a preprocessor check rather than a comment asking the next reader to be careful, so reordering the file breaks the build instead of quietly testing the plain entry point. The `StreamCaptureBeginSpt` roundtrip asserts the `_spt` begin id and the plain end id together, which is what proves the staging did what it claims. The patterns are repeated bytes (`0x41`, `0x42`, `0x44`, `0x45`, `0x46`, `0x4A`, `0x4C`, `0x5A`) for the reason #8840 documented: an expected value has to decode above `atol/(1-rtol)` in every encoding the D2H validator tries before an all-zero replay buffer is reported as a mismatch. `Unit_HRR_StreamCaptureBeginSpt` previously used `0x33`, which decodes to 4.17e-8 as f32 and so passed against zeros under the default tolerance; the roundtrips now also pin `HIP_HRR_D2H_EXACT=1`. `hrr_run_exact_roundtrip()` takes an optional list of expected API ids, and `hrr_require_recorded_apis()` generalises the id assertions #8840 added to `Unit_HRR_MemsetSptRoundtrip`. ## Issue Tracking JIRA ID: AIRUNTIME-2665 ## Test Plan Build `libamdhip64`, `hrr-playback` and the `HrrTest` Catch2 target on a gfx950 host, then run the full HRR unit suite as a regression pass. ## Test Result Verified on a gfx950 host at this revision: the full HRR unit suite passes, both the GPU-free pass (19/19) and the GPU pass (52/52), with `Unit_HRR_MemcpyPeerRoundtrip` skipped as it needs a second GPU. All nine per-thread roundtrips pass, the four added here and the five that already existed. The four workloads whose `_spt` blocks were removed (`MemsetExtra`, `MemcpyExtra`, `StreamAdvanced2`, `HostRegLaunch`) still pass. The build is itself part of the evidence for the macro staging: the guard before the memcpy workload fails the build if `hipMemcpy` is no longer redirected there, and the guard after the `#undef` fails it if the redirect survived, so a clean compile pins the macro state at each stage. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <Daniel.Olsson@amd.com> Co-authored-by: Daniel Olsson <113136742+danielbolsson@users.noreply.github.com>

Copilot AI lite review requested due to automatic review settings August 18, 2026 11:09

andraghetti requested review from a team as code owners August 18, 2026 11:09

github-actions Bot added the project: hip-tests label Aug 18, 2026

systems-assistant Bot added the organization: ROCm label Aug 18, 2026

Copilot started reviewing on behalf of andraghetti August 18, 2026 11:10

View session

Copilot AI reviewed Aug 18, 2026

View reviewed changes

Comment thread projects/hip-tests/catch/unit/hrr/hrr_spt_workload_test.cc Resolved

Comment thread projects/hip-tests/catch/config/configs/unit/hrr.yaml Resolved

andraghetti force-pushed the users/andraghetti/hrr-spt-test-conversion branch from 843110f to 75134e7 Compare August 18, 2026 12:23

andraghetti changed the title ~~test(hrr): [AIRUNTIME-2665] convert _spt test calls to compile-option pattern~~ test(hrr): [AIRUNTIME-2665] reach the remaining _spt APIs via the compile option Aug 18, 2026

gandryey approved these changes Aug 19, 2026

View reviewed changes

danielbolsson and others added 8 commits August 21, 2026 02:16

Merge remote-tracking branch 'origin/develop' into users/andraghetti/…

3f2e2a6

…hrr-spt-test-conversion

Merge branch 'develop' into users/andraghetti/hrr-spt-test-conversion

b91cf7e

Merge branch 'develop' into users/andraghetti/hrr-spt-test-conversion

208340c

test(hrr): disable Unit_HRR_*Spt cases on amd_windows/linux/wsl

789d606

Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>

Merge branch 'develop' into users/andraghetti/hrr-spt-test-conversion

9b87884

Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>

Merge branch 'develop' into users/andraghetti/hrr-spt-test-conversion

39e614c

Merge branch 'develop' into users/andraghetti/hrr-spt-test-conversion

df76e48

github-actions Bot added the project: hrr label Sep 9, 2026

danielbolsson merged commit a2efae8 into develop Sep 9, 2026
26 of 27 checks passed

danielbolsson deleted the users/andraghetti/hrr-spt-test-conversion branch September 9, 2026 08:51

Navigation Menu

Repository navigation

test(hrr): [AIRUNTIME-2665] reach the remaining _spt APIs via the compile option - #10306

test(hrr): [AIRUNTIME-2665] reach the remaining _spt APIs via the compile option#10306
danielbolsson merged 9 commits into
developfrom
users/andraghetti/hrr-spt-test-conversion

andraghetti commented Aug 18, 2026 •

edited

Loading

Uh oh!

therock-pr-bot Bot commented Aug 18, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

systems-assistant Bot commented Sep 10, 2026

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

andraghetti commented Aug 18, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Aug 18, 2026

✅ All Policy Checks Passed

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

systems-assistant Bot commented Sep 10, 2026

TheRock Submodule Bump Activity

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

andraghetti commented Aug 18, 2026 •

edited

Loading