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

projects/clr/hipamd/src/hip_code_object.cpp

@@ -708,4 +708,31 @@ void StatCO::ForEachFatBinaryBlob(void (*cb)(const void*)) const {
     cb(data);
   }
 }
+
+// ================================================================================================
+void StatCO::ForEachGlobalVar(void (*cb)(const void*, const char*, size_t,
+                                         const void*)) {
+  std::scoped_lock lock(sclock_);
+  // Resolved here rather than by the caller: this sweep runs from HRR's
+  // capture init, which is inside hip::init(), and re-entering the runtime
+  // through hipGetSymbolAddress from there deadlocks on the init once-flag.
+  // No device is current that early, and GetGlobalVar asserts rather than
+  // erroring on a negative device id, so the address is reported as
+  // unresolved instead.
+  // No device is current this early in init, and GetGlobalVar asserts rather
+  // than erroring on a negative id, so fall back to device 0 — which is the
+  // device a recording resolves its globals on anyway.
+  int device_id = ihipGetDevice();
+  if (device_id < 0) device_id = 0;
+  const bool resolvable = static_cast<size_t>(device_id) < g_devices.size();
+  for (const auto& [host_var, var] : vars_) {
+    if (var == nullptr) continue;
+    hipDeviceptr_t dev_ptr = nullptr;
+    size_t sym_size = 0;
+    if (resolvable &&
+        GetGlobalVar(host_var, device_id, &dev_ptr, &sym_size) != hipSuccess)
+      dev_ptr = nullptr;
+    cb(host_var, var->GetName().c_str(), var->GetSize(), dev_ptr);
+  }
+}
 }  // namespace hip

projects/clr/hipamd/src/hip_code_object.hpp

@@ -193,6 +193,12 @@ class StatCO : public CodeObject {
   // Iterate all registered fat binary data pointers — for HRR capture post-registration sweep.
   void ForEachFatBinaryBlob(void (*cb)(const void*)) const;
 
+  // Iterate all registered __device__ globals as (host shadow address, symbol
+  // name, size, device address) — the same post-registration sweep for HRR,
+  // which otherwise never sees __hipRegisterVar because it fires at
+  // static-init time. The device address is null if it cannot be resolved yet.
+  void ForEachGlobalVar(void (*cb)(const void*, const char*, size_t, const void*));
+
  private:
   mutable std::recursive_mutex sclock_;    //!< Guards Static Code object
   const PlatformState& owner_;             //!< Reference to owning PlatformState

projects/clr/hipamd/src/hrr/hip_capture_writer.cpp

@@ -737,14 +737,15 @@ void checkpoint() {
 void mark_incomplete(const char* reason) {
   // Record once; the loud, AMD_LOG_LEVEL-routed message is emitted by the caller
   // (e.g. serialize_kernel_launch) which has the relevant context. Here we only
-  // need the durable flag and a single stderr breadcrumb so a bare run still
-  // surfaces it.
+  // need the durable flag and a single breadcrumb so a bare run still surfaces
+  // it. Error level: the archive is not faithful and replay must not treat it
+  // as one. log_printf appends its own newline, so the format string omits it.
   if (!g_capture_incomplete.exchange(true, std::memory_order_relaxed)) {
-    fprintf(stderr,
-            "[HRR capture] Archive marked INCOMPLETE: %s. The clean-shutdown "
-            "trailer will be omitted and manifest.complete=false so replay "
-            "cannot treat this capture as faithful.\n",
-            reason ? reason : "(unspecified)");
+    LogPrintfError(
+        "[HRR capture] Archive marked INCOMPLETE: %s. The clean-shutdown "
+        "trailer will be omitted and manifest.complete=false so replay "
+        "cannot treat this capture as faithful",
+        reason ? reason : "(unspecified)");
   }
 }
 

projects/hrr/DESIGN.md

@@ -131,7 +131,7 @@ precisely the ones left without a trailer and absent from the root index, while
 the parent that exited cleanly needs no repair. Sub-archives that already carry
 a clean trailer are skipped without being read.
 
-### Archive Format (v5)
+### Archive Format (v6)
 ```
 capture.hrr/
   manifest.json      { version, capture_mode, owner_pid, processes[] }
@@ -242,7 +242,7 @@ projects/hrr/                     — standalone HRR project (portable layer)
                                     Not used by capture — it is the format an
                                     out-of-tree producer writes and playback reads
   playback/
-    hrr_reader.h/.cpp             — archive loader; record framing
+    hrr_reader.h/.cpp             — archive loader, v6 format; record framing
                                     (read_raw_record / open_record_stream) shared
                                     with the region sidecars
     hrr_region_map.h/.cpp         — region timeline: merge, cursor, live block set,
@@ -321,13 +321,13 @@ The generator classifies each API:
 Generated capture shims for manual APIs are pass-throughs (no `write_event()`).
 When adding HIP API support, update this script to classify the API in the appropriate capture and playback policy sets. APIs requiring non-trivial serialization or replay belong in `MANUAL_CAPTURE_APIS` and/or `MANUAL_PLAYBACK_APIS`; intentionally unsupported replay APIs belong in `NOOP_PLAYBACK_APIS`.
 
-## Archive Format (v5)
+## Archive Format (v6)
 
 Single-authority definition in `hrr_api_args.h` (auto-generated):
 
 ```
 HRR_MAGIC   = 0x52524845  ("HRRE")
-HRR_VERSION = 5
+HRR_VERSION = 6
 ```
 
 Version history, so an archive written by an older runtime can be placed:
@@ -357,7 +357,7 @@ Version history, so an archive written by an older runtime can be placed:
 <output_dir>/
   manifest.json      { version, capture_mode, owner_pid, processes[] }
                      (version here is the manifest schema = 1, distinct from the
-                      events.bin HRR_VERSION = 5)
+                      events.bin HRR_VERSION = 6)
   pid-<pid>/
     manifest.json      { pid, parent_pid, complete, event_count, blob_count }
     writer_state.json  checkpoint cursor (next_seq, event/blob counts, events file
@@ -974,24 +974,20 @@ There is no practical fix: capturing general CPU synchronisation would require
 instrumentation of the application binary, OS synchronisation primitives, or a
 language runtime — all of which are out of scope for a HIP API trace.
 
-### `__hipRegisterVar` / `__hipRegisterManagedVar` — Not Replayed
+### `__hipRegisterManagedVar` — Not Replayed
 
 These compiler-generated functions populate the runtime's internal host-symbol →
 device-symbol table used by `hipMemcpyToSymbol`, `hipMemcpyFromSymbol`,
 `hipGetSymbolAddress`, and `hipGetSymbolSize`.
 
-The capture shims record events for `__hipRegisterVar` and `__hipRegisterManagedVar`,
-but the symbol name is stored only as a raw `uint64_t` host pointer (a capture-time
-address, useless at replay) rather than the name string, and the playback shims are
-no-ops. The host-symbol → device-symbol table is therefore never rebuilt.
+`__hipRegisterVar` is handled — see "Symbol Registry" below: the name string is recorded
+and replay rebuilds the table through `hipModuleGetGlobal`, so the `__device__` global
+families replay. `__hipRegisterManagedVar` is not: a `__managed__` global is never
+registered in a way the sweep sees, so its symbol never reaches `symbol_map`.
 
-**Impact:** workloads that use `hipMemcpyToSymbol` / `hipMemcpyFromSymbol` will fail
-at replay because the symbol table is never populated. Kernel launches and all
-pointer-based memcpy APIs are unaffected.
-
-**Workaround:** none currently. A fix would require recording the variable name and
-host pointer, then at playback calling `hipModuleGetGlobal` on the loaded module to
-resolve the device address and rebuilding the symbol table entry.
+**Impact:** a workload whose `hipMemcpy*Symbol` subject is a `__managed__` rather than a
+`__device__` global fails loudly at replay, naming the symbol it could not resolve.
+Kernel launches and all pointer-based memcpy APIs are unaffected.
 
 ### `co_hash` Kernel → Code-Object Binding
 
@@ -1284,31 +1280,47 @@ Other HIP features with similar complications include device-side enqueue, coope
 groups with CPU-side barriers, and persistent kernels with CPU-side steering. These are
 not currently handled.
 
-### Explicit Graph Construction — Not Supported (fails loudly)
-
-Only the **stream-capture** graph chain is supported (`hipStreamBeginCapture` →
-`hipStreamEndCapture` → `hipGraphInstantiate` → `hipGraphLaunch`). Graphs built
-explicitly through the node API are not — and replay no longer silently produces an
-empty graph with skipped launches (finding H1). The failure is split into two levels:
-
-- **Hard fail-loud at instantiate.** `hipGraphInstantiate` / `hipGraphInstantiateWithFlags`
-  (manual handlers) return `hipErrorNotSupported` — which the dispatcher treats as
-  fatal — when the graph handle is absent from `graph_map`. A miss can only happen for
-  a node-API-built graph, since the stream-capture chain always records its graph via
-  `hipStreamEndCapture`. This is the point that actually matters: a non-replayable
-  graph can never be launched as an empty graph.
-- **Attributable warning at construction (non-fatal).** `hipGraphCreate` /
-  `hipGraphClone` / every `hipGraphAdd*Node` / `hipGraphInstantiateWithParams` are
-  `ERROR_STUB_PLAYBACK_APIS`: they log a loud, per-API "explicit (node-API) graph
-  construction is NOT supported" message and return `hipSuccess`. They are deliberately
-  non-fatal so a program that merely *creates/clones* a graph (or exercises these APIs
-  for coverage) without instantiating it in an unsupported way still replays; if such a
-  graph is later instantiated, the instantiate gate above fires.
-- Node-parameter struct pointers are still stored as raw `uint64_t` without
-  dereferencing (full node-API replay remains future work).
-
-The supported stream-capture path is unaffected: it never emits `hipGraphCreate` /
-`hipGraphAdd*Node` events, and its graph is in `graph_map` so instantiate succeeds.
+### Explicit Graph Construction — Supported, With One Refusal
+
+Both graph chains replay: **stream capture** (`hipStreamBeginCapture` →
+`hipStreamEndCapture` → `hipGraphInstantiate` → `hipGraphLaunch`) and **explicit
+node-API construction** (`hipGraphCreate` → `hipGraphAdd*Node` → `hipGraphInstantiate`).
+A node-API graph used to replay as an empty graph with its launches skipped (finding
+H1); it is now rebuilt node by node:
+
+- `graph_node_map` in `PlaybackContext` maps each recorded `hipGraphNode_t` to the node
+  replay created, which is what lets dependency arrays, `*NodeGetParams` and the
+  `*NodeSetParams` / `*ExecNodeSetParams` mutation family refer to earlier nodes.
+- Every `hipGraphAdd*Node` construction API has a real handler. Node-parameter structs
+  are deep-copied at capture through `DEREF_FIELDS` (kernel node params reuse the
+  kernel-argument encoder), and at replay the pointers, streams, events and symbols
+  inside them are translated before the node is created.
+- The mutation spellings are real handlers rather than no-ops, so a graph whose
+  parameters change between launches replays with the parameters the recording used
+  instead of every launch repeating the first one.
+
+The fail-loud gate at instantiate stays as the backstop. `hipGraphInstantiate` returns
+`hipErrorNotSupported` when the graph handle is absent from `graph_map`, or when the
+graph contains a node HRR cannot reproduce — a host-function node, since host callbacks
+are a declared scope exclusion. A graph missing work the recording had is refused rather
+than launched.
+
+### Symbol Registry — `__device__` Globals Only
+
+Replay loads code objects with `hipModuleLoadData` and so has no host shadow variables,
+while the `hipMemcpy*Symbol`, `hipGetSymbolAddress/Size` and graph symbol-node families
+all take a host shadow address as their subject. Capture bridges the gap by sweeping the
+registered `__device__` globals (`StatCO::ForEachGlobalVar`) and recording each one's
+*name* alongside the recorded host address; replay resolves the name through
+`hipModuleGetGlobal` and keeps the result in `symbol_map`. The graph symbol nodes are
+then spelled as 1D memcpy nodes against the resolved device address, which is
+semantically the same node without needing a host shadow to point at.
+
+The sweep runs from `hip_capture_init()`, which is itself called from `hip::init()`, so
+it must not call back into a public HIP entry point: doing so re-enters the
+initialisation once-guard and deadlocks. Resolution therefore happens inside
+`ForEachGlobalVar` while it already holds the code-object lock, falling back to device 0
+when no device is current yet.
 
 ### Wire-Format Size Limits
 
@@ -1319,7 +1331,7 @@ The event wire format (finding H5):
   shrinking `reserved` to 2 bytes), so kernel launches with large serialized payloads
   (many args / long mangled names / large by-value structs) up to ~4 GiB are recorded
   normally instead of being dropped at 65535 bytes. This is the change that bumped
-  `HRR_VERSION` to 4; the current version is 5, see Archive Format above. The
+  `HRR_VERSION` to 4; the current version is 6, see Archive Format above. The
   writer's single-record buffer path now writes any oversized record straight through.
 - **Per-argument size limit (64 KiB) now fails loudly.** Each kernel arg's size is still
   a `uint16_t`. A by-value struct argument ≥ 64 KiB cannot be represented, so the launch
@@ -1398,15 +1410,27 @@ resolver scans **all** loaded modules and binds (and caches) the first same-name
 function from any code object, with no identity check — potentially binding the wrong
 kernel.
 
-### Host Callbacks and Struct-Pointer Inputs — Not Captured
+### Struct-Pointer Inputs — Declared, Not Silent
+
+The generator no longer drops a pointed-to struct by default. `DEREF_FIELDS` declares,
+per API and parameter, how much the pointer points at — a fixed struct, `count`
+elements of a fixed size, a NUL-terminated string, or an array whose elements are
+themselves device addresses — and the generator emits both the capture-side copy and
+the playback-side reconstruct-and-pass. The recorded capture-time address is carried for
+comparison only; the value handed back to the runtime is always a local of the pointee's
+own type.
+
+What is left is the honest residue: an API whose pointer argument has no `DEREF_FIELDS`
+entry still records only the address. That is now visible rather than silent — the
+manifest classifies it as payload loss and `derive_manifest.py` fails the build if the
+count drifts from its recorded baseline, so a newly-captured struct-input API cannot
+lose its payload unnoticed.
+
+### Recorded Failures Are Not Replay Failures
 
-- `hipStreamAddCallback` / `hipLaunchHostFunc` record the callback pointer as 0; the
-  host-side work and any device-staging side effects it performs are never captured or
-  replayed. This is distinct from the app-thread CPU-sync limitation above.
-- By default the code generator stores an input struct-pointer parameter as a raw
-  `uint64_t` without serialising the pointed-to struct; only hand-written manual shims
-  copy struct bytes. A newly-captured, non-manual struct-input API therefore loses its
-  payload silently.
+A recorded call that *failed* is reproduced, not repaired: when the archived return is
+non-zero and replay reproduces the same error, that is the faithful outcome and the
+dispatcher reports it as such instead of as a handler failure.
 
 ## Relationship to Original HRR Code