andraghetti · Oct 16, 2026 · Oct 18, 2026 · Oct 18, 2026 · Oct 19, 2026 · Oct 19, 2026
@@ -725,10 +725,10 @@ hipError_t StatCO::InitManagedVarDevicePtr(int deviceId) {
 }
 
 // ================================================================================================
-void StatCO::ForEachFatBinaryBlob(void (*cb)(const void*)) const {
+void StatCO::ForEachFatBinary(void (*cb)(FatBinaryInfo*)) {
   std::scoped_lock lock(sclock_);
-  for (const auto& [data, _] : modules_) {
-    cb(data);
+  for (const auto& [_, fat_binary_info] : modules_) {
+    cb(fat_binary_info);
   }
 }
 

@@ -197,8 +197,9 @@ class StatCO : public CodeObject {
   // pointer is resolvable through MemObjMap before any HIP API can observe it.
   hipError_t EnsureAllManagedVarsAreMapped();
 
-  // Iterate all registered fat binary data pointers — for HRR capture post-registration sweep.
-  void ForEachFatBinaryBlob(void (*cb)(const void*)) const;
+  // Iterate all registered fat binaries' FatBinaryInfo, null until HIP digests
+  // it; for HRR capture post-registration sweep.
+  void ForEachFatBinary(void (*cb)(FatBinaryInfo*));
 
   // Iterate all registered __device__ globals as (host shadow address, symbol
   // name, size, device address) — the same post-registration sweep for HRR,

@@ -1024,9 +1024,8 @@ static uint64_t file_bytes(uint64_t len) {
 // still opens an archive of its own, held to the same limit.
 // HIP_HRR_CAPTURE_MAX_BLOB_MB caps one blob: a larger copy is not stored, its
 // event records no blob, the archive is marked incomplete and capture goes on.
-// Code objects and the fat binaries HIP registers are not held to it; a
-// hipLinkAddData image is. 0, the default, sets no limit. Both apply with or
-// without the reserve.
+// Code objects are not held to it; a hipLinkAddData image is. 0, the default,
+// sets no limit. Both apply with or without the reserve.
 //
 // A write over the archive limit is refused and capture stopped under the same
 // hold of g_file_mu, which flush() takes to decide on the trailer: an archive
@@ -1811,13 +1810,13 @@ class BlobClaim {
 // write_blob
 // ---------------------------------------------------------------------------
 
-Hash128 write_blob(const void* data, size_t len, bool capped) {
+Hash128 write_blob(const void* data, size_t len) {
   reopen_after_fork();
   {
     std::lock_guard<std::mutex> lk(g_file_mu);
     if (g_events_fd < 0 || g_events_finalized) return {};  // not open, or past the trailer
   }
-  if (capped && g_max_blob_bytes != 0 && len > g_max_blob_bytes) {
+  if (g_max_blob_bytes != 0 && len > g_max_blob_bytes) {
     skip_oversized_blob(len);
     return {};
   }

@@ -91,12 +91,10 @@ void note_unreplayable(const char* api, const char* reason);
 
 // Write a buffer as a content-addressed blob. Returns hash, or {} when the
 // writer is not open, flush() has written the trailer, or capture has stopped
-// for lack of space or at the size limit. A `capped` buffer over
-// HIP_HRR_CAPTURE_MAX_BLOB_MB is not written either: {} is returned and the
-// capture marked incomplete. Fat binaries are not capped: replay loads them to
-// run their kernels and resolve their device globals.
+// for lack of space or at the size limit. A buffer over HIP_HRR_CAPTURE_MAX_BLOB_MB
+// is not written either: {} is returned and the capture marked incomplete.
 // Thread-safe. Skips the write when the same bytes are already in this archive.
-Hash128 write_blob(const void* data, size_t len, bool capped = true);
+Hash128 write_blob(const void* data, size_t len);
 
 // Write a code object (.hsaco) blob. Returns hash, or {} as write_blob() does.
 Hash128 write_code_object(const void* image, size_t image_size);

@@ -52,7 +52,7 @@ through one of these shims, which records a fixed-size `hrr_event_header` (32 by
 type, sequence ID, timestamp, thread ID) followed by a variable-length payload of
 serialised arguments into a streaming `events.bin` file.
 
-A subset of APIs requires hand-written shims (`MANUAL_CAPTURE_APIS`): kernel launches (4 variants), host-to-device memcpy (4 variants), module load (3 variants), `__hipRegisterFatBinary`, host memory registration (`hipHostRegister`, `hipHostUnregister`), all `hipMemcpy3D` variants (4), array creation, VMM, stream/memory attribute APIs, and others where struct pointer serialisation requires custom handling. All others are code-generated by `gen_hrr_api_args.py` from `hip_api_trace.hpp`. Fat binaries registered before shim install (C++ static-init) are recovered retroactively via `ForEachFatBinaryBlob` at init.
+A subset of APIs requires hand-written shims (`MANUAL_CAPTURE_APIS`): kernel launches (4 variants), host-to-device memcpy (4 variants), module load (3 variants), `__hipRegisterFatBinary`, host memory registration (`hipHostRegister`, `hipHostUnregister`), all `hipMemcpy3D` variants (4), array creation, VMM, stream/memory attribute APIs, and others where struct pointer serialisation requires custom handling. All others are code-generated by `gen_hrr_api_args.py` from `hip_api_trace.hpp`. Fat binaries registered before shim install (C++ static-init) are recovered retroactively via `ForEachFatBinary` at init.
 
 H2D memcpy captures the host source buffer as a content-addressed blob (FNV-1a-128 hash). D2H memcpy captures the expected output blob after a forced stream sync. Kernel argument pointers are detected via `kernel->signature()` (in-tree access to `amd::KernelSignature`) — no ELF parsing needed.
 
@@ -189,7 +189,7 @@ capture.hrr/
     manifest.json      { pid, parent_pid, complete, event_count, blob_count }
     active             names this process instance: present while its capture is on (see Transport)
     blobs/<2hex>/      content-addressed host buffers keyed by FNV-1a-128 hash
-    code_objects/      .hsaco ELFs (unused in current fat-binary path)
+    code_objects/      .hsaco ELFs: loaded modules and the code object HIP takes from each fat binary
     regions/*.hrrr     external region annotations (optional; written by producers
                        outside the runtime, never by capture — see below)
   pid-<pid>.prev-<n>/  an earlier process's archive with the same pid, moved
@@ -508,7 +508,7 @@ manually: `const uint8_t* p = pl + sizeof(hrr_event_header);`
 
 | API | Extra fields |
 |-----|-------------|
-| `__hipRegisterFatBinary` | `blob_hash_lo u64`, `blob_hash_hi u64`, `blob_size u64` |
+| `__hipRegisterFatBinary` | `blob_hash_lo u64`, `blob_hash_hi u64`, `blob_size u64` (0 when the hash names the code object HIP loaded) |
 | `hipModuleLoadData`, `hipModuleLoadDataEx`, `hipModuleLoad` | `co_hash_lo u64`, `co_hash_hi u64`, `module_id u32` |
 | `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync` | `blob_hash_lo u64`, `blob_hash_hi u64` |
 | `hipHostRegister` | `blob_hash_lo u64`, `blob_hash_hi u64` |
@@ -794,9 +794,15 @@ not its type, so the type is not compared.
 
 `__hipRegisterFatBinary` fires at C++ static-init time before shims are installed.
 At init, all already-registered fat binaries are recovered retroactively via
-`ForEachFatBinaryBlob`. At playback, each captured blob is loaded with
-`hipModuleLoadData`; the resulting `hipModule_t` is stored in `module_map` keyed by
-hash and scanned for kernel name lookups.
+`ForEachFatBinary`. Capture does not store the bundle: its size could only come
+from the header, which nothing bounds. It records the code object HIP extracted
+from the bundle for the current device, bounded by the ELF's own headers and by the
+memory that can be read, since the size HIP knows is the one the bundle entry declares
+and the headers can declare as much; one whose headers reach further is not recorded.
+The register event names that code object (blob hash set, blob size 0). Replay pre-loads every code object,
+so kernels and `__device__` globals resolve as they do for module loads. A register
+event with a blob size carries the fat binary itself as a blob, which playback loads
+with `hipModuleLoadData`.
 
 ## Multi-Threaded Capture
 
@@ -1212,7 +1218,7 @@ kernel.
 `__hipRegisterFatBinary` fires at C++ static-initialisation time, before
 `hip_capture_init()` installs shims. All fat binaries registered before the shim is
 active are missed by the normal capture path and recovered retroactively via
-`PlatformState::Instance().StatCO().ForEachFatBinaryBlob(...)` at init time.
+`PlatformState::Instance().StatCO().ForEachFatBinary(...)` at init time.
 
 If an application loads additional shared libraries _after_ HIP init that register
 their own fat binaries, those will be captured normally through the shim. The edge

@@ -57,11 +57,10 @@ not make.
 | Host buffers copied to or from the device, byte for byte, unencrypted. For a serving stack this includes prompts and generated output as token ids, which the model's public tokeniser reverses, and the model weights | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) |
 | Host memory registered with `hipHostRegister`: the whole range, byte for byte, as it was at registration, whether or not it is ever copied | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) |
 | Kernel argument bytes | `events.bin` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) |
-| Code objects the process loaded (the application's own GPU binary code) | `code_objects/`, `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) |
+| Code objects the process loaded (the application's own GPU binary code); of a fat binary, the code object HIP loads, which its `__hipRegisterFatBinary` record names with a blob size of 0. A record with a blob size names a whole bundle in `blobs/`, with targets HIP never loads | `code_objects/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880); [AIRUNTIME-2858](https://amd-hub.atlassian.net/browse/AIRUNTIME-2858): of a fat binary, only the code object HIP loads |
 | Linker input buffers passed to `hipLinkAddData`, every byte, with their names, including inputs the call rejected and inputs never linked or loaded. A SPIR-V or LLVM bitcode input keeps its compiler identification and, built with `-g`, its source and build directories; capture does not scrub them | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) |
-| Whole fat binary bundles, including targets HIP never loads | `blobs/` | contents | [AIRUNTIME-2858](https://amd-hub.atlassian.net/browse/AIRUNTIME-2858): only the code objects HIP loads |
-| Debug and comment sections of code objects: absolute build paths, source file names, the exact compiler build | `code_objects/`, `blobs/` | contents | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): zeroed in code objects and uncompressed fat binaries; kept in compressed (`CCOB`) fat binaries |
-| Strings in the code the GPU loads, such as the absolute source path a device `assert()` prints | `code_objects/`, `blobs/` | contents | — kept, the kernel uses them; `verify` counts code objects holding an absolute path |
+| Debug and comment sections of code objects: absolute build paths, source file names, the exact compiler build | `code_objects/` | contents | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): zeroed in every code object stored, including the one HIP takes from a compressed (`CCOB`) fat binary. A compressed bundle that a record with a blob size names in `blobs/` keeps them |
+| Strings in the code the GPU loads, such as the absolute source path a device `assert()` prints | `code_objects/` | contents | — kept, the kernel uses them; `verify` counts code objects holding an absolute path |
 | Framework allocator layout, when a region producer ran beside capture | `regions/` | both | — kept, written only by an opt-in producer |
 | Marker that capture is running: on Linux the boot id and the process start time in clock ticks since boot; on Windows the process creation date and time. A normal exit removes it; a process that crashes or is killed leaves it | `active` | both | — kept, region producers read it |
 

@@ -94,7 +94,7 @@ User-facing capture, replay, and validation knobs. Implementation details can be
 | `HIP_HRR_DEBUG_ARGS` | off | Dump every captured kernel arg to the log (debug / provenance). Only a Debug build of the HIP runtime (no `NDEBUG`) reads it; a release build ignores it. Either way capture says on stderr whether it reads it, whatever the log level |
 | `HIP_HRR_HOST_SNAPSHOTS` | on | Record the pinned host memory each kernel reads before the launch, so replay can write it back. `HIP_HRR_HOST_SNAPSHOT_MAX_MB` (default `64`) caps the allocation size. The archive then holds that host memory (see DESIGN.md, Pinned Host Snapshots) |
 | `HIP_HRR_CAPTURE_MAX_MB` | 0 (no limit) | Largest size, in MiB, of each process's archive (`pid-<pid>/`), manifests aside. The write that would pass it is not made: capture stops with a `[HRR capture] Capture stopped` message on stderr, the archive is marked incomplete, and the application runs on. A process it forks afterwards still gets an archive of its own |
-| `HIP_HRR_CAPTURE_MAX_BLOB_MB` | 0 (no limit) | Largest copy, in MiB, stored as a blob. A larger one is recorded without its contents, the archive is marked incomplete with one message on stderr, and capture goes on. Replay refuses by name a host-to-device copy recorded so, whether `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync`, `hipMemcpyWithStream`, a 2D or 3D copy or a to-symbol graph copy, a `hipHostRegister`, and a kernel launch outside a stream capture whose pinned host snapshot was left out. Code objects and the fat binaries HIP registers are always stored; a `hipLinkAddData` image is not, and replay refuses its link |
+| `HIP_HRR_CAPTURE_MAX_BLOB_MB` | 0 (no limit) | Largest copy, in MiB, stored as a blob. A larger one is recorded without its contents, the archive is marked incomplete with one message on stderr, and capture goes on. Replay refuses by name a host-to-device copy recorded so, whether `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync`, `hipMemcpyWithStream`, a 2D or 3D copy or a to-symbol graph copy, a `hipHostRegister`, and a kernel launch outside a stream capture whose pinned host snapshot was left out. Code objects are always stored; a `hipLinkAddData` image is not, and replay refuses its link |
 
 On Linux the archive is readable only by the user who captured it (directories 0700, files 0600). What an archive contains, which of it each capture mode records, and why replaying an archive means running its GPU code: [PRIVACY.md](PRIVACY.md). If the archive cannot be set up, capture is disabled with a `[HRR capture] Capture disabled` message on stderr and the application runs normally.
 
@@ -248,11 +248,12 @@ capture.hrr/
   still resolve to the code object they came from. Their `.comment` and
   `.debug*` sections, which hold the compiler build and the build paths of
   debug information, are zeroed and retyped `SHT_NOBITS` before they are
-  stored, as are those of uncompressed fat binaries, so a debugger on a
-  replay sees no debug information. What the GPU loads is not touched:
-  kernel names stay, and so does a string a kernel keeps, such as the
-  source path a device `assert()` prints. Compressed (`CCOB`) fat binaries
-  and `hipLinkAddData` inputs are stored as they were passed.
+  stored, so a debugger on a replay sees no debug information. Of a fat
+  binary, compressed or not, only the code object HIP loads is stored, here
+  and scrubbed the same way. What the GPU loads is not touched: kernel
+  names stay, and so does a string a kernel keeps, such as the source path
+  a device `assert()` prints. `hipLinkAddData` inputs are stored as they
+  were passed.
 - **manifest.json** — counts, versions and device properties; the GPU's PCI
   address, UUID and LUID are written blank
 - **Complete: NO** — original run crashed before clean shutdown; reader still recovers complete events

@@ -210,7 +210,7 @@ typedef struct {
 typedef struct {
     hrr_event_header hdr;
     uint64_t ret;
-    uint64_t blob_hash_lo;  /* FNV-1a-128 lo of fat binary */
+    uint64_t blob_hash_lo;  /* FNV-1a-128 lo of fat binary, or of its code object if blob_size is 0 */
     uint64_t blob_hash_hi;  /* FNV-1a-128 hi */
     uint64_t blob_size;  /* fat binary byte count */
 } hrr_args___hipRegisterFatBinary;

@@ -1759,6 +1759,7 @@ hipError_t playback___hipRegisterFatBinary(PlaybackContext& ctx,
     uint64_t blob_hash_hi = a->blob_hash_hi;
 
     if (!blob_hash_lo && !blob_hash_hi) return hipSuccess;  // no blob — skip
+    // No blob: the hash names a code object, which run_pass() pre-loads.
     if (!a->blob_size) return hipSuccess;
 
     std::string hex = hrr::hash_hex(blob_hash_lo, blob_hash_hi);