Skip to content

feat(hrr): record pointed-to payloads and replay the handler families - #10641

Merged
ammallya merged 36 commits into
developfrom
users/danielbolsson/hrr-api-payloads
Sep 28, 2026
Merged

ammallya merged 36 commits into
developfrom
users/danielbolsson/hrr-api-payloads

Conversation

@danielbolsson

@danielbolsson danielbolsson commented Aug 24, 2026 •

Copy link
Copy Markdown
Contributor

AIRUNTIME-2703, AIRUNTIME-2704 · sixth of nine in the stack rooted at #10221 · base is #10637

What this changes

An API whose argument is a pointer to a struct used to record only the address. Replay had nothing to reconstruct from, so 44 such APIs were either hand-written shims or no-op stubs, and the ones left as stubs lost their payload silently.

DEREF_FIELDS declares, per API and parameter, what the pointer points at — a fixed struct, count elements of a fixed size, a NUL-terminated string, or an array of device addresses — and the generator emits both the capture-side copy and the playback-side reconstruct-and-pass. The recorded capture-time address is kept for comparison only; the value handed to the runtime is always a local of the pointee's own type. Output handle parameters get the same treatment through _OUT_BUFFER_PARAMS, which is what stops an export writing into a stale capture-time address.

_OUT_BUFFER_PARAMS also has to cover the callee-filled buffer that is not a handle. hipDeviceGetLuid takes a char*, which the pointee rule would otherwise size as one character for the runtime to memcpy eight bytes into. It gets an explicit 8. Only the Windows path writes, so the overflow is unreachable on Linux and no Linux-only test run can find it — worth an eye on review rather than on CI.

On top of that mechanism, the handler families that needed more than a declaration: graphs rebuilt node by node against graph_node_map, the symbol registry resolving __device__ globals through hipModuleGetGlobal, ipc_handle_map, the driver ctx_map and link_state_map, and the memcpy2D/3D and batch-memop families. Archive format goes to v6.

A call that failed at capture and fails the same way at replay is fidelity, not a replay error, and the dispatcher now reports it as such instead of aborting.

Reviewing this

Hand-written gen_hrr_api_args.py, hip_capture.cpp, hip_playback.cpp/.h, hip_code_object.cpp/.hpp, DESIGN.md
Generated hrr_api_args.h, hip_capture_generated.cpp, hip_playback_generated.cpp — run gen_hrr_api_args.py, confirm an empty diff
Expectations delta hrr_api_matrix_expectations.h — the rows this PR reclassifies, relative to #10637's baseline

One deliberate stub

hipMemPoolImportFromShareableHandle becomes a no-op here. Its export side now genuinely mints an fd, so the recorded handle names something else in the replaying process and faults inside the pool import. The PR above replaces the stub with a refusal that names the API and the reason.

Validation

Full HRR unit suite on gfx950, GPU 7: 86/86, both the GPU-free pass (26/26) and the GPU pass. All four runnable api-matrix tiers pass.

Made with Cursor

@LeoZhao-BJ

LeoZhao-BJ commented Sep 11, 2026 •

Copy link
Copy Markdown
Contributor

I think we should brainstorm a more generic approach for handling HRR payload capture/replay, instead of continuing to add or customize shims API by API.

Two directions seem worth exploring:

  1. Keep a declarative per-API spec, but centralize the generated capture/replay logic

    Instead of hand-writing or expanding bespoke shim logic for each affected API, we could describe pointer payloads in a compact metadata table, then let common
    capture/replay helpers handle the actual deep-copy and reconstruction.

    For example, the generator could emit specs such as:

    hipGraphAddMemcpyNode:
    pDependencies: array<hipGraphNode_t>, count=numDependencies
    pCopyParams: struct

    hipMemcpyParam2DAsync:
    pCopy: struct<hip_Memcpy2D>

    The generated shim would then call shared helpers like:

    capture_deep_fields(api_id, args, field_specs)
    replay_rebuild_fields(ctx, args, field_specs)

    This would keep the per-API knowledge declarative, while moving the risky behavior — null handling, truncation policy, host/device pointer translation,
    lifetime, and refusal logic — into one place.

  2. Move deep-copied payloads into a generic TLV-style event tail

    Instead of growing each hrr_args_* struct with fields like foo_bytes, foo_present, and foo_n, we could encode deep-copied payloads as generic typed field
    records appended after the fixed event payload.

    Each captured field could carry:

    field_id
    field_kind
    element_size
    original_count
    recorded_count
    byte_size
    bytes[]

    This would let replay distinguish:

    • fully captured payloads;
    • truncated payloads;
    • absent/null payloads;
    • host buffers;
    • device pointer arrays;
    • unsupported/unreplayable pointer kinds.

    It would also make policy enforcement generic. For example, if original_count > recorded_count, replay should fail loudly or mark the graph/event incomplete,
    rather than silently replaying only the captured prefix and treating the partial operation as faithful.

The broader goal would be to classify pointer parameters by semantic kind, not just by C type:

StructByPointer
ArrayByCount
String
DevicePointer
DevicePointerArray
HostInputBuffer
HostOutputBuffer
OutputHandle
CallbackFunction
OSHandle
JitOptionValueArray
Unreplayable

Then each semantic kind has a common capture/replay policy. Most APIs could be handled by generated metadata and shared helpers, while only genuinely semantic families — kernel launches, graph lifecycle, IPC handle remapping, JIT option interpretation, symbol registration — would need hand-written replay logic.

This should reduce the need to add bespoke shims for every API while also making the behavior safer and more consistent across the replay stack.

danielbolsson and others added 5 commits September 14, 2026 15:54
…lsson/hrr-api-payloads

Semantic conflict in the generator and its generated playback.

Both sides insert code immediately before _playback_arg and #10641 also
changed that function's signature to take the ApiEntry. The two evolutions
are orthogonal and both are wanted, so both bodies are kept and the new
signature wins:

* #10641's argument derivation (out buffers, by-value structs carried as
  inline bytes, dereferenced pointees, live-ctx and post-call tables).
* #10632's _PLAYBACK_ARG_BRIDGES and validate_playback_signatures(), which
  refuse to emit a playback cast that disagrees with the public API the
  call is passed to.

They meet on hipMemAdvise: #10641 leaves it on the scalar fallback, which
still emits the dispatch-table `(int)a->device` that the public API does not
take at HIP_FORCE_API_VERSION >= 800. The validation catches it and the
bridge resolves it, so the generated handler now carries the version guard
and #10641's recorded-error block together.

_apply_arg_bridge substitutes by position, so it now asserts that the
generated call has one argument per dispatch-table parameter rather than
trusting that invariant across future edits.

The generated files are the generator's output, not a hand merge: they were
regenerated and verified idempotent. That corrects one pre-existing staleness
in #10641's committed playback — hipChooseDevice is cast to
hipDeviceProp_tR0600, the spelling its own dispatch-table typedef has always
used. #10641's own generator on its own tree emits the same thing, so this is
the committed file catching up, not a behaviour change (the two names are the
same type by typedef).

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…ix' into users/danielbolsson/hrr-api-payloads
…ix' into users/danielbolsson/hrr-api-payloads
…ix' into users/danielbolsson/hrr-api-payloads
@danielbolsson

Copy link
Copy Markdown
Contributor Author

I think we should brainstorm a more generic approach for handling HRR payload capture/replay, instead of continuing to add or customize shims API by API.

Two directions seem worth exploring:

  1. Keep a declarative per-API spec, but centralize the generated capture/replay logic
    Instead of hand-writing or expanding bespoke shim logic for each affected API, we could describe pointer payloads in a compact metadata table, then let common
    capture/replay helpers handle the actual deep-copy and reconstruction.
    For example, the generator could emit specs such as:
    hipGraphAddMemcpyNode:
    pDependencies: array<hipGraphNode_t>, count=numDependencies
    pCopyParams: struct
    hipMemcpyParam2DAsync:
    pCopy: struct<hip_Memcpy2D>
    The generated shim would then call shared helpers like:
    capture_deep_fields(api_id, args, field_specs)
    replay_rebuild_fields(ctx, args, field_specs)
    This would keep the per-API knowledge declarative, while moving the risky behavior — null handling, truncation policy, host/device pointer translation,
    lifetime, and refusal logic — into one place.

This is largely built already. gen_hrr_api_args.py has a Deref class that is exactly the declarative per-API spec as you propose. Your proposal is basically a rename of something already there.

Your strongest point here is the truncation one, and it is a real gap today. Let's adress that!

  1. Move deep-copied payloads into a generic TLV-style event tail
    Instead of growing each hrr_args_* struct with fields like foo_bytes, foo_present, and foo_n, we could encode deep-copied payloads as generic typed field
    records appended after the fixed event payload.

This is an interesting architectural change proposal. I suggest that we first land the current PR stack and discuss future design directions from there.

Raw fprintf(stderr, ...) is not permitted in clr: it bypasses the
AMD_LOG_LEVEL machinery and writes unconditionally from inside a library
that other processes embed.

Convert the two capture diagnostics the generator emits, and the
archive-incomplete notice in the writer, to LogPrintfWarning and
LogPrintfError from utils/debug.hpp. The generated file now includes that
header; the writer already did. Trailing newlines are dropped from the
format strings because amd::log_printf appends its own.

The 32 converted call sites in hip_capture_generated.cpp come from two
emit sites in gen_hrr_api_args.py. Regenerating is a fixed point: two
consecutive runs produce byte-identical output.

Deliberately unchanged:
  - hip_playback_generated.cpp keeps fprintf. It builds into hrr-playback,
    a standalone executable that links the public HIP package and not CLR
    internals, so utils/debug.hpp is unavailable to it by design.
  - The manifest.json writes in hip_capture_writer.cpp are file I/O to an
    open FILE*, not logging.

Known tradeoff: AMD_LOG_LEVEL defaults to 0, so these messages are now
silent unless logging is raised. They report a truncated or incomplete
capture, which is exactly when a user most needs to see something, so this
may warrant a mechanism that surfaces integrity warnings regardless of log
level.

@LeoZhao-BJ LeoZhao-BJ 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.

let's land this PR first and may have epic to optimize architecture later

danielbolsson and others added 12 commits September 17, 2026 06:55
hip_capture_generated.cpp does not compile: every LogPrintfWarning in it
fails with "AMD_LOG_LEVEL was not declared in this scope".

debug.hpp's ClPrint macro reads the AMD_LOG_LEVEL and AMD_LOG_MASK flag
variables but does not declare them; they come from utils/flags.hpp,
whose defaults in turn need top.hpp for the unit suffixes (Mi) and
typedefs (uint) they are written with. hip_capture_writer.cpp gets all
three transitively through os/os.hpp, which is why it builds and the
generated file does not.

Add top.hpp and utils/flags.hpp ahead of utils/debug.hpp in the generated
capture preamble.

Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>
…ix' into users/danielbolsson/hrr-api-payloads
…ix' into users/danielbolsson/hrr-api-payloads
Raw fprintf(stderr, ...) is not permitted in clr: it bypasses the
AMD_LOG_LEVEL machinery and writes unconditionally from inside a library
that other processes embed. mark_incomplete() is the last such call site
in this file, so it is routed through LogPrintfError from utils/debug.hpp,
which the file already includes. Error level rather than warning: once
this fires the archive is not faithful and replay must not treat it as
one. The trailing newline is dropped because amd::log_printf appends its
own.

The conversion was reverted on #10634 (cd4d649) so that PR's existing
approval stays valid. That revert is an ancestor of this branch, so
cascading it up the stack silently restored the raw fprintf here with no
conflict. This commit re-applies the conversion explicitly, so the fix
still reaches develop with this PR while #10634 keeps its approved shape.

The converted text is byte-identical to the form already carried by the
rest of the stack, which is what keeps the remaining merges conflict-free.

Signed-off-by: Daniel Olsson <Daniel.Olsson@amd.com>
…n/hrr-api-payloads

# Conflicts:
#	projects/clr/hipamd/src/hrr/hip_capture.cpp

Co-authored-by: Daniel Olsson <danielbolsson@users.noreply.github.com>
…n/hrr-api-payloads

Co-authored-by: Daniel Olsson <danielbolsson@users.noreply.github.com>
# Conflicts:
#	projects/hrr/tools/gen_hrr_api_args.py
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

7 participants