From 582005f049f89f5ed2546041c689c30e8a2ca994 Mon Sep 17 00:00:00 2001 From: andraghetti Date: Fri, 16 Oct 2026 13:35:24 +0200 Subject: [PATCH 1/8] feat(hrr): strip build info from stored code objects, blank GPU identifiers A capture meant to be shared carried the build tree and the machine with it: code objects keep the compiler build in .comment and, when built with debug information, source file names and absolute build paths in .debug*; the manifest named the GPU by PCI address, UUID and LUID. Replay needs none of it. Capture now zeroes the contents of every .comment and .debug* section before it stores a code object, both the device images under code_objects/ and each entry of an uncompressed fat binary under blobs/. Zeroing in place keeps sizes and offsets, so a bundle's entry table stays valid and the image loads as before; the stored bytes are hashed after scrubbing, so the archive stays content-addressed. Compressed (CCOB) fat binaries are stored as they were. The manifest keeps the pci, uuid and luid keys and writes them blank. inspect_archive.py verify reports, per process, which hardware identifiers are still set and how many code objects still hold build information, so the person approving a transfer can see it. Jira: AIRUNTIME-2881 --- projects/clr/hipamd/src/hrr/CMakeLists.txt | 1 + projects/clr/hipamd/src/hrr/hip_capture.cpp | 14 +- .../hipamd/src/hrr/hip_capture_metadata.cpp | 26 +--- .../clr/hipamd/src/hrr/hip_capture_scrub.cpp | 119 ++++++++++++++ .../clr/hipamd/src/hrr/hip_capture_scrub.h | 24 +++ .../clr/hipamd/src/hrr/hip_capture_writer.cpp | 11 +- projects/hrr/README.md | 6 +- projects/hrr/skills/enable-recording/SKILL.md | 8 +- .../scripts/inspect_archive.py | 98 ++++++++++++ .../tests/test_inspect_archive.py | 78 ++++++++++ .../integration/hrr_fatbin_header_test.cc | 145 +++++++++++++++++- .../tests/integration/hrr_roundtrip_test.cc | 13 +- 12 files changed, 508 insertions(+), 35 deletions(-) create mode 100644 projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp create mode 100644 projects/clr/hipamd/src/hrr/hip_capture_scrub.h diff --git a/projects/clr/hipamd/src/hrr/CMakeLists.txt b/projects/clr/hipamd/src/hrr/CMakeLists.txt index 6cb64bf620f9..2119e402e061 100644 --- a/projects/clr/hipamd/src/hrr/CMakeLists.txt +++ b/projects/clr/hipamd/src/hrr/CMakeLists.txt @@ -21,6 +21,7 @@ endif() target_sources(amdhip64 PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/hip_capture.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hip_capture_metadata.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/hip_capture_scrub.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hip_capture_writer.cpp ${CMAKE_CURRENT_SOURCE_DIR}/hip_capture_generated.cpp ) diff --git a/projects/clr/hipamd/src/hrr/hip_capture.cpp b/projects/clr/hipamd/src/hrr/hip_capture.cpp index 57cf4dd96d76..27d1565af35b 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture.cpp @@ -31,6 +31,7 @@ #include "hip_capture.h" #include "hip_capture_metadata.h" +#include "hip_capture_scrub.h" #include "hip_capture_writer.h" // hrr_api_args.h — for hrr_args_* struct types and hrr_api_id_t enum @@ -1436,6 +1437,15 @@ static size_t compute_bundle_size(const void* blob) { return static_cast(end); } +// Stores a fat binary with the build paths and compiler build of its code +// objects zeroed; the application's own copy is left alone. +static hrr_cap::Hash128 write_fat_binary_blob(const void* blob, size_t blob_size) { + std::vector scrubbed(static_cast(blob), + static_cast(blob) + blob_size); + hrr_cap::scrub::scrub_code_object(scrubbed.data(), scrubbed.size()); + return hrr_cap::writer::write_blob(scrubbed.data(), scrubbed.size()); +} + void** capture___hipRegisterFatBinary(const void* data) { void** r = g_real_compiler_table.__hipRegisterFatBinary_fn(data); // Shim is only installed when capture is active — no hip_capture_enabled() check needed. @@ -1464,7 +1474,7 @@ void** capture___hipRegisterFatBinary(const void* data) { a.ret = reinterpret_cast(r); a.blob_size = static_cast(blob_size); if (blob && blob_size > 0) { - auto h = hrr_cap::writer::write_blob(blob, blob_size); + auto h = write_fat_binary_blob(blob, blob_size); a.blob_hash_lo = h.lo; a.blob_hash_hi = h.hi; } @@ -2622,7 +2632,7 @@ static void record_fat_binary_blob(const void* blob_ptr) { hrr_args___hipRegisterFatBinary a{}; a.ret = 0; // handle not meaningful at init time a.blob_size = static_cast(blob_size); - auto h = hrr_cap::writer::write_blob(blob_ptr, blob_size); + auto h = write_fat_binary_blob(blob_ptr, blob_size); a.blob_hash_lo = h.lo; a.blob_hash_hi = h.hi; hrr_cap::writer::write_event_raw(HRR_API_HIPREGISTERFATBINARY, &a.hdr, sizeof(a)); diff --git a/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp b/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp index a6678e5ea4ef..e2f1c566817b 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_metadata.cpp @@ -88,22 +88,6 @@ std::string version_string_from_int(int version) { return std::to_string(major) + "." + std::to_string(minor) + "." + std::to_string(patch); } -std::string bytes_to_hex(const char* bytes, size_t len) { - static constexpr char kHex[] = "0123456789abcdef"; - std::string result; - result.resize(len * 2); - for (size_t i = 0; i < len; ++i) { - const auto v = static_cast(bytes[i]); - result[2 * i] = kHex[v >> 4]; - result[2 * i + 1] = kHex[v & 0xf]; - } - return result; -} - -std::string uuid_to_hex(const hipUUID& uuid) { - return bytes_to_hex(uuid.bytes, sizeof(uuid.bytes)); -} - void append_prop_fields(std::ostringstream& os, const hipDeviceProp_t& prop) { os << " \"properties\": {\n" << " \"name\": " << quote(bounded_string(prop.name, sizeof(prop.name))) << ",\n" @@ -123,11 +107,11 @@ void append_prop_fields(std::ostringstream& os, const hipDeviceProp_t& prop) { << " \"compute_mode\": " << prop.computeMode << ",\n" << " \"compute_capability\": " << quote(std::to_string(prop.major) + "." + std::to_string(prop.minor)) << ",\n" - << " \"pci\": " << quote(std::to_string(prop.pciDomainID) + ":" + - std::to_string(prop.pciBusID) + ":" + - std::to_string(prop.pciDeviceID)) << ",\n" - << " \"uuid\": " << quote(uuid_to_hex(prop.uuid)) << ",\n" - << " \"luid\": " << quote(bytes_to_hex(prop.luid, sizeof(prop.luid))) << ",\n" + // Left blank: they name the individual accelerator, and with it the machine + // that captured, and replay needs none of them. + << " \"pci\": \"\",\n" + << " \"uuid\": \"\",\n" + << " \"luid\": \"\",\n" << " \"luid_device_node_mask\": " << prop.luidDeviceNodeMask << "\n" << " }"; } diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp new file mode 100644 index 000000000000..0996001e07f3 --- /dev/null +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp @@ -0,0 +1,119 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +#include "hip_capture_scrub.h" + +#include + +// Field offsets are spelled out rather than taken from , which Windows +// does not have and whose Elf64_* typedefs collide with CLR's own ELF headers. +// The image is the application's memory, so every read is bounds-checked and +// copied out with memcpy, whatever its alignment. + +namespace hrr_cap { +namespace scrub { + +namespace { + +constexpr size_t kElfHeaderSize = 64; +constexpr size_t kSectionHeaderSize = 64; +constexpr uint8_t kElfClass64 = 2; +constexpr uint8_t kElfDataLsb = 1; +constexpr uint32_t kShtNobits = 8; +constexpr uint64_t kShfAlloc = 0x2; + +constexpr char kBundleMagic[] = "__CLANG_OFFLOAD_BUNDLE__"; +constexpr size_t kBundleMagicLen = sizeof(kBundleMagic) - 1; +constexpr size_t kBundleEntryHeaderSize = 3 * sizeof(uint64_t); + +template +T read(const uint8_t* p, size_t at) { + T v; + std::memcpy(&v, p + at, sizeof(v)); + return v; +} + +// [offset, offset + size) lies inside [0, limit). +bool in_range(uint64_t offset, uint64_t size, size_t limit) { + return offset <= limit && size <= limit - offset; +} + +bool has_prefix(const char* s, size_t len, const char* prefix) { + const size_t n = std::strlen(prefix); + return len >= n && std::memcmp(s, prefix, n) == 0; +} + +bool is_scrubbed_section(const char* name, size_t len) { + return (len == 8 && std::memcmp(name, ".comment", 8) == 0) || has_prefix(name, len, ".debug"); +} + +size_t scrub_elf(uint8_t* image, size_t size) { + if (size < kElfHeaderSize || std::memcmp(image, "\x7f" "ELF", 4) != 0 || + image[4] != kElfClass64 || image[5] != kElfDataLsb) + return 0; + const uint64_t shoff = read(image, 0x28); + const uint16_t shentsize = read(image, 0x3a); + const uint16_t shnum = read(image, 0x3c); + const uint16_t shstrndx = read(image, 0x3e); + // shnum == 0 also covers extended section numbering, which code objects never need. + if (shentsize != kSectionHeaderSize || shnum == 0 || shstrndx >= shnum || + !in_range(shoff, uint64_t{shnum} * kSectionHeaderSize, size)) + return 0; + + const uint8_t* shstr_hdr = image + shoff + size_t{shstrndx} * kSectionHeaderSize; + const uint64_t names_off = read(shstr_hdr, 24); + const uint64_t names_size = read(shstr_hdr, 32); + if (!in_range(names_off, names_size, size)) return 0; + const char* names = reinterpret_cast(image + names_off); + + size_t zeroed = 0; + for (uint16_t i = 0; i < shnum; ++i) { + const uint8_t* sh = image + shoff + size_t{i} * kSectionHeaderSize; + const uint32_t name = read(sh, 0); + const uint32_t type = read(sh, 4); + const uint64_t flags = read(sh, 8); + const uint64_t offset = read(sh, 24); + const uint64_t sec_size = read(sh, 32); + if (name >= names_size || type == kShtNobits || (flags & kShfAlloc) != 0) continue; + const size_t name_len = strnlen(names + name, names_size - name); + if (!is_scrubbed_section(names + name, name_len)) continue; + if (sec_size == 0 || !in_range(offset, sec_size, size)) continue; + std::memset(image + offset, 0, sec_size); + ++zeroed; + } + return zeroed; +} + +size_t scrub_bundle(uint8_t* image, size_t size) { + size_t pos = kBundleMagicLen; + if (size < pos + sizeof(uint64_t)) return 0; + const uint64_t count = read(image, pos); + pos += sizeof(uint64_t); + size_t zeroed = 0; + for (uint64_t i = 0; i < count; ++i) { + if (size - pos < kBundleEntryHeaderSize) break; + const uint64_t offset = read(image, pos); + const uint64_t entry_size = read(image, pos + 8); + const uint64_t id_size = read(image, pos + 16); + pos += kBundleEntryHeaderSize; + if (id_size > size - pos) break; + pos += id_size; + if (in_range(offset, entry_size, size)) zeroed += scrub_elf(image + offset, entry_size); + } + return zeroed; +} + +} // namespace + +size_t scrub_code_object(uint8_t* image, size_t size) { + if (!image) return 0; + if (size >= kBundleMagicLen && std::memcmp(image, kBundleMagic, kBundleMagicLen) == 0) + return scrub_bundle(image, size); + return scrub_elf(image, size); +} + +} // namespace scrub +} // namespace hrr_cap diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.h b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h new file mode 100644 index 000000000000..84baf850c942 --- /dev/null +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h @@ -0,0 +1,24 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ +#pragma once + +#include +#include + +namespace hrr_cap { +namespace scrub { + +// Zeroes, in place, the contents of the .comment and .debug* sections of a code +// object about to be stored: an ELF64 image, or every entry of an uncompressed +// clang offload bundle. Those sections carry the compiler build, absolute build +// paths and source file names; the GPU never loads them, so replay is unchanged. +// Sizes and offsets stay as they are, so a bundle's entry table stays valid. +// A compressed bundle, or anything that does not parse, is left as it is. +// Returns how many sections were zeroed. +size_t scrub_code_object(uint8_t* image, size_t size); + +} // namespace scrub +} // namespace hrr_cap diff --git a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp index 6a4cff183160..9aa0982992a7 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp @@ -39,6 +39,7 @@ #include "hip_capture_writer.h" #include "hip_capture.h" #include "hip_capture_metadata.h" +#include "hip_capture_scrub.h" #include "os/os.hpp" // amd::Os::timeNanos() #include "utils/debug.hpp" // LogPrintfError, LogPrintfWarning, LogPrintfInfo @@ -1638,7 +1639,13 @@ Hash128 write_code_object(const void* image, size_t image_size) { if (g_events_fd < 0 || g_events_finalized) return {}; // not open, or past the trailer } - Hash128 h = hash_buffer(image, image_size); + // The stored bytes, and so the hash, are the image without its build paths + // and compiler build (see hip_capture_scrub.h). + std::vector scrubbed(static_cast(image), + static_cast(image) + image_size); + scrub::scrub_code_object(scrubbed.data(), scrubbed.size()); + + Hash128 h = hash_buffer(scrubbed.data(), image_size); char hex[33]; hash_hex(h, hex); std::string key = std::string("co:") + hex; // namespace to match playback load_code_object key @@ -1658,7 +1665,7 @@ Hash128 write_code_object(const void* image, size_t image_size) { ClaimedWrite claim(key); std::string path = g_output_dir + "/code_objects/" + hex + ".hsaco"; - const bool ok = atomic_write_file(path, image, image_size); + const bool ok = atomic_write_file(path, scrubbed.data(), image_size); if (!ok) LogPrintfWarning("[HRR capture] Failed to write code object %s", hex); claim.done(ok); return h; diff --git a/projects/hrr/README.md b/projects/hrr/README.md index 02bad2355acf..ff3ea0edafd4 100644 --- a/projects/hrr/README.md +++ b/projects/hrr/README.md @@ -235,7 +235,11 @@ capture.hrr/ - **code_objects/** — code objects, content-addressed by hash (device images extracted by the runtime, for `hipModuleLoad` too); a launch records the hash so kernels sharing a name (Triton emits many `triton_`) - still resolve to the code object they came from + still resolve to the code object they came from. Their `.comment` and + `.debug*` sections, which hold the compiler build and build paths, are + zeroed before they are stored, as are those of uncompressed fat binaries +- **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 Code objects read back from the runtime must start with `\x7fELF`, diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index 33721c99a87d..808554476e77 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -124,19 +124,19 @@ An archive with several `pid-*` directories is one archive. Send the whole direc ## 4. Share it, knowing what is in it -**There is no scrubbing, anonymisation or randomisation option today.** Say that plainly before anyone asks whether the archive can be sanitised, because the answer shapes what they do next and there is no flag that changes it. +**There is no scrubbing, anonymisation or randomisation option for the data today.** Say that plainly before anyone asks whether the archive can be sanitised, because the answer shapes what they do next and there is no flag that changes it. Capture does drop two things on its own, with no flag: the build paths and compiler build inside the code objects, and the GPU's hardware identifiers. What the archive holds, per process, and it is worth reading out to whoever has to approve the transfer: - **`events.bin`**: every successful HIP call the process made, with its arguments: sizes, device addresses, stream handles, kernel names. - **`blobs/`**: the host buffers those calls carried, byte for byte and in the clear. For a serving stack that means prompts and generated text; for training, the data and the weights that were copied to the device. -- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. -- **`manifest.json`**: process and parent process ids, event and blob counts, the clean-shutdown flag, the HIP and comgr versions that captured, and the device's name, architecture, PCI location and UUID. +- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. Their `.comment` and `.debug*` sections are zeroed, here and in the fat binaries under `blobs/`, so the compiler build, source file names and build paths are gone, except in compressed (`CCOB`) fat binaries, which are stored as the application had them. Kernel names stay. +- **`manifest.json`**: process and parent process ids, event and blob counts, the clean-shutdown flag, the HIP and comgr versions that captured, and the device's name, architecture, memory size and clocks. Its PCI location, UUID and LUID are written blank. - **`regions/*.hrrr`**: the framework allocator's own layout, when a producer ran alongside capture. Apart from the few exceptions in step 2, failed HIP calls are not recorded, and nothing that never crossed a HIP call is. The archive is exactly as sensitive as the data the workload handled, so where it may go has to be agreed before it is uploaded. -`verify --json` gives each process's counts and sizes in a form you can attach to the request; the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. +`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`, `build_info_files`, and `unchecked_fat_binaries` for the compressed ones it cannot read); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. If the content is the obstacle, three things exist and none of them is scrubbing. Re-capture the same failure driven by synthetic input, if it still reproduces. Replay the archive yourself and send only the finding, which is a few lines and carries no buffers. Or agree the handling of a production-data transfer before uploading, because that is what this is. diff --git a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py index 840814609155..541b1574a492 100755 --- a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py +++ b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py @@ -73,6 +73,14 @@ # in playback/hrr_reader.cpp), however large the file is. MAX_RECORD_BYTES = 64 * 1024 * 1024 +# Capture zeroes these ELF sections in every code object it stores, since they +# hold the compiler build, source file names and build paths; and it writes the +# manifest's hardware identifiers blank. `verify` reports what is left of both. +# See projects/clr/hipamd/src/hrr/hip_capture_scrub.h. +BUNDLE_MAGIC = b"__CLANG_OFFLOAD_BUNDLE__" +COMPRESSED_BUNDLE_MAGIC = b"CCOB" +HARDWARE_ID_FIELDS = ("pci", "uuid", "luid") + @dataclass class ProcessArchive: @@ -92,6 +100,12 @@ class ProcessArchive: hip_runtime_version: str | None = None format_version: int | None = None holds_an_event: bool = False + # Manifest fields that name the individual GPU and are not blank. + hardware_ids: list[str] = field(default_factory=list) + # Code objects and fat binaries still holding a .comment or .debug* section. + build_info_files: int = 0 + # Compressed fat binaries, which cannot be checked without decompressing. + unchecked_fat_binaries: int = 0 @property def recorded(self) -> bool: @@ -166,6 +180,75 @@ def _dir_stats(path: Path) -> tuple[int, int]: return files, total +def _elf_has_build_info(data: bytes) -> bool: + """Does this ELF64 image hold a .comment or .debug* section that is not all zero?""" + if len(data) < 64 or data[:4] != b"\x7fELF" or data[4] != 2 or data[5] != 1: + return False + shoff = struct.unpack_from("= shnum or shoff + shnum * 64 > len(data): + return False + names = struct.unpack_from("= 0 else len(data)] + if label != b".comment" and not label.startswith(b".debug"): + continue + if kind == 8 or flags & 0x2: # SHT_NOBITS, SHF_ALLOC + continue + if data[offset : offset + size].strip(b"\0"): + return True + return False + + +def _code_object_has_build_info(data: bytes) -> bool: + """The same, for an ELF or for any entry of an uncompressed offload bundle.""" + if not data.startswith(BUNDLE_MAGIC): + return _elf_has_build_info(data) + pos = len(BUNDLE_MAGIC) + if len(data) < pos + 8: + return False + (count,) = struct.unpack_from(" None: + """Count the stored code objects that still carry build information. + + Every file under `code_objects/` is one; under `blobs/` only the fat + binaries are, and the first bytes tell them from buffer contents without + reading the buffers. + """ + candidates = sorted((pid_dir / "code_objects").glob("*")) + sorted( + (pid_dir / "blobs").rglob("*") + ) + for path in candidates: + try: + if not path.is_file(): + continue + with path.open("rb") as handle: + head = handle.read(len(BUNDLE_MAGIC)) + if head.startswith(COMPRESSED_BUNDLE_MAGIC): + proc.unchecked_fat_binaries += 1 + continue + if path.parent.name != "code_objects" and head != BUNDLE_MAGIC: + continue + data = head + handle.read() + except OSError: + continue + if _code_object_has_build_info(data): + proc.build_info_files += 1 + + def _archive_format_version(events: Path) -> int | None: """The format version written in the first 8 bytes of `events.bin`.""" try: @@ -246,6 +329,7 @@ def _load_process(pid_dir: Path) -> ProcessArchive: proc.holds_an_event = _holds_an_event(events) proc.blob_files, _ = _dir_stats(pid_dir / "blobs") proc.code_object_files, _ = _dir_stats(pid_dir / "code_objects") + _scan_code_objects(pid_dir, proc) manifest, error = _read_manifest(pid_dir / "manifest.json") proc.manifest_error = error @@ -280,6 +364,12 @@ def typed(source: dict, key: str, kind: type): runtime = metadata.get("runtime") if isinstance(metadata, dict) else None if isinstance(runtime, dict): proc.hip_runtime_version = typed(runtime, "hip_runtime_version", str) + devices = metadata.get("devices") if isinstance(metadata, dict) else None + for device in devices if isinstance(devices, list) else []: + props = device.get("properties") if isinstance(device, dict) else None + for key in HARDWARE_ID_FIELDS: + if isinstance(props, dict) and props.get(key) and key not in proc.hardware_ids: + proc.hardware_ids.append(key) if proc.pid is None and pid_dir.name.startswith("pid-"): suffix = pid_dir.name[len("pid-") :] @@ -703,6 +793,11 @@ def render(report: ArchiveReport) -> str: ) if proc.parent_pid: lines.append(f" parent pid: {proc.parent_pid}") + hardware = ", ".join(proc.hardware_ids) if proc.hardware_ids else "blank" + build_info = f"build info in {proc.build_info_files} code objects" + if proc.unchecked_fat_binaries: + build_info += f", {proc.unchecked_fat_binaries} compressed fat binaries not checked" + lines.append(f" hardware identifiers: {hardware} {build_info}") incomplete = [p for p in report.processes if p.complete is False] if incomplete: @@ -767,6 +862,9 @@ def to_dict(report: ArchiveReport) -> dict: "blob_files": p.blob_files, "code_object_files": p.code_object_files, "finalized": p.finalized, + "hardware_ids": p.hardware_ids, + "build_info_files": p.build_info_files, + "unchecked_fat_binaries": p.unchecked_fat_binaries, } for p in report.processes ], diff --git a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py index 23547d0eab93..05f7bbbe7ec7 100644 --- a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py +++ b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py @@ -562,6 +562,84 @@ def test_json_output_is_machine_readable(tmp_path, capsys): assert payload["processes"][0]["pid"] == 99 +def make_elf(sections: dict[str, bytes]) -> bytes: + """A minimal ELF64 little-endian image holding `sections`, then .shstrtab.""" + names = b"\0" + offsets = {} + for name in [*sections, ".shstrtab"]: + offsets[name] = len(names) + names += name.encode() + b"\0" + body = b"\0" * 64 + headers = [struct.pack(" bytes: + """An uncompressed offload bundle: an empty host entry, then `elf`.""" + host, target = b"host-x86_64-unknown-linux-gnu-", b"hipv4-amdgcn-amd-amdhsa--gfx950" + table = inspect_archive.BUNDLE_MAGIC + struct.pack(" +#include #include #include #include @@ -203,14 +208,15 @@ void** register_fat_binary(const void* bundle) { return __hipRegisterFatBinary(new Wrapper{kHipfMagic, 1, bundle, nullptr}); } -std::vector compile_code_object() { +std::vector compile_code_object(std::vector options = {}) { static const char* src = R"( extern "C" __global__ void hrr_fatbin_probe(int* out) { out[0] = 1; } )"; hiprtcProgram prog = nullptr; HRR_HIPRTC_CHECK(hiprtcCreateProgram(&prog, src, "hrr_fatbin_probe.hip", 0, nullptr, nullptr)); - if (hiprtcCompileProgram(prog, 0, nullptr) != HIPRTC_SUCCESS) { + if (hiprtcCompileProgram(prog, static_cast(options.size()), options.data()) != + HIPRTC_SUCCESS) { size_t log_sz = 0; (void)hiprtcGetProgramLogSize(prog, &log_sz); std::string log(log_sz, '\0'); @@ -387,6 +393,141 @@ HRR_TEST_CASE(Unit_HRR_FatBinHeaderChecks) { CHECK(out.find("HIP did not register it") != std::string::npos); } +// =========================================================================== +// Code objects are stored without their build paths and compiler build. +// =========================================================================== +TEST_CASE("Unit_HRR_CodeObjectScrub_Direct", "[.][hrr-direct]") { + HRR_HIP_CHECK(hipSetDevice(0)); + hipDeviceProp_t prop{}; + HRR_HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + const std::string target = std::string(kTargetPrefix) + prop.gcnArchName; + // Debug information is what puts the source file name and build directory + // into a code object. + const std::vector co = compile_code_object({"-g"}); + + // Stored as a fat binary blob. + REQUIRE(register_fat_binary(pin(valid_bundle(target, co))) != nullptr); + + // Stored under code_objects/, and launched so replay runs the stored copy. + hipModule_t mod = nullptr; + HRR_HIP_CHECK(hipModuleLoadData(&mod, co.data())); + hipFunction_t fn = nullptr; + HRR_HIP_CHECK(hipModuleGetFunction(&fn, mod, "hrr_fatbin_probe")); + int* d = nullptr; + HRR_HIP_CHECK(hipMalloc(&d, sizeof(int))); + HRR_HIP_CHECK(hipMemset(d, 0, sizeof(int))); + void* args[] = {&d}; + HRR_HIP_CHECK(hipModuleLaunchKernel(fn, 1, 1, 1, 1, 1, 1, 0, nullptr, args, nullptr)); + int h = 0; + HRR_HIP_CHECK(hipMemcpy(&h, d, sizeof(int), hipMemcpyDeviceToHost)); + REQUIRE(h == 1); + HRR_HIP_CHECK(hipFree(d)); + HRR_HIP_CHECK(hipModuleUnload(mod)); +} + +namespace { +struct ScrubCount { + size_t elves = 0; + size_t sections = 0; // .comment and .debug* sections with contents + size_t left = 0; // of those, still holding a non-zero byte +}; + +// Walks the section headers of an ELF64 image independently of capture's own +// parser and counts the sections capture has to zero. +void count_scrub_sections(const uint8_t* p, size_t size, ScrubCount& c) { + auto u16 = [&](size_t at) { uint16_t v; std::memcpy(&v, p + at, 2); return v; }; + auto u32 = [&](size_t at) { uint32_t v; std::memcpy(&v, p + at, 4); return v; }; + auto u64 = [&](size_t at) { uint64_t v; std::memcpy(&v, p + at, 8); return v; }; + REQUIRE(size >= 64); + REQUIRE(std::memcmp(p, "\x7f" "ELF", 4) == 0); + ++c.elves; + const uint64_t shoff = u64(0x28); + const uint16_t shnum = u16(0x3c); + const uint16_t shstrndx = u16(0x3e); + REQUIRE(u16(0x3a) == 64); + REQUIRE(shoff + uint64_t{shnum} * 64 <= size); + REQUIRE(shstrndx < shnum); + const uint64_t names = u64(shoff + shstrndx * 64 + 24); + for (uint16_t i = 0; i < shnum; ++i) { + const size_t sh = shoff + size_t{i} * 64; + REQUIRE(names + u32(sh) < size); + const std::string name(reinterpret_cast(p + names + u32(sh))); + if (name != ".comment" && name.rfind(".debug", 0) != 0) continue; + const uint64_t off = u64(sh + 24), len = u64(sh + 32); + if (u32(sh + 4) == 8 /*SHT_NOBITS*/ || len == 0) continue; + REQUIRE(off + len <= size); + INFO("Section " << name); + ++c.sections; + for (uint64_t b = 0; b < len; ++b) { + if (p[off + b] != 0) { + ++c.left; + break; + } + } + } +} + +bool contains(const std::vector& bytes, const std::string& needle) { + return std::search(bytes.begin(), bytes.end(), needle.begin(), needle.end()) != bytes.end(); +} +} // namespace + +HRR_TEST_CASE(Unit_HRR_CodeObjectScrub) { + ScopedDir cap(fs::temp_directory_path() / "hrr_code_object_scrub.hrr"); + // Replay checks the kernel's result, so the scrubbed copy loads and runs. + hrr_run_roundtrip("Unit_HRR_CodeObjectScrub_Direct", cap.path); + + hrr::Archive arc; + REQUIRE(hrr::load_archive(cap.path.string(), arc)); + const fs::path archive = hrr_single_process_archive(cap.path); + + ScrubCount module_count; + size_t module_files = 0; + for (const auto& entry : fs::directory_iterator(archive / "code_objects")) { + INFO("Code object " << entry.path()); + const std::string text = read_text_file(entry.path()); + const std::vector image(text.begin(), text.end()); + ++module_files; + count_scrub_sections(image.data(), image.size(), module_count); + CHECK_FALSE(contains(image, "clang version")); + CHECK_FALSE(contains(image, "hrr_fatbin_probe.hip")); + } + REQUIRE(module_files >= 1); + + // The bundle's entries are walked from its own table, which zeroing in place + // left as it was. + ScrubCount bundle_count; + for (const auto& ev : arc.events) { + if (ev.header().event_type != HRR_API_HIPREGISTERFATBINARY) continue; + const auto* a = + reinterpret_cast(ev.raw_payload.data()); + if (a->blob_hash_lo == 0 && a->blob_hash_hi == 0) continue; + std::vector blob; + REQUIRE(hrr::read_blob(arc, a->blob_hash_lo, a->blob_hash_hi, blob)); + if (blob.size() < kTableStart || std::memcmp(blob.data(), kBundleMagic, kBundleMagicLen) != 0) + continue; // a compressed bundle from the test binary itself + size_t pos = kTableStart; + for (uint64_t i = 0, n = read_u64(blob, kBundleMagicLen); i < n; ++i) { + const uint64_t off = read_u64(blob, pos), len = read_u64(blob, pos + 8); + pos += kEntryHeaderSize + read_u64(blob, pos + 16); + if (len >= 4 && std::memcmp(blob.data() + off, "\x7f" "ELF", 4) == 0) + count_scrub_sections(blob.data() + off, len, bundle_count); + } + CHECK_FALSE(contains(blob, "clang version")); + CHECK_FALSE(contains(blob, "hrr_fatbin_probe.hip")); + } + + INFO("Module code objects: " << module_count.elves << " ELF, " << module_count.sections + << " sections, " << module_count.left << " not zeroed"); + INFO("Fat binary code objects: " << bundle_count.elves << " ELF, " << bundle_count.sections + << " sections, " << bundle_count.left << " not zeroed"); + // The probe was built with -g, so both copies had something to zero. + CHECK(module_count.sections >= 2); + CHECK(bundle_count.sections >= 2); + CHECK(module_count.left == 0); + CHECK(bundle_count.left == 0); +} + #endif // HRR_TEST_EXE /** diff --git a/projects/hrr/tests/integration/hrr_roundtrip_test.cc b/projects/hrr/tests/integration/hrr_roundtrip_test.cc index 92f06a26e904..db166824fefa 100644 --- a/projects/hrr/tests/integration/hrr_roundtrip_test.cc +++ b/projects/hrr/tests/integration/hrr_roundtrip_test.cc @@ -1060,6 +1060,7 @@ HRR_TEST_CASE(Unit_HRR_MetadataManifest) { const std::string compute_capability = json_string_value(metadata, "compute_capability"); const std::string pci = json_string_value(metadata, "pci"); const std::string uuid = json_string_value(metadata, "uuid"); + const std::string luid = json_string_value(metadata, "luid"); CHECK(ordinal >= 0); CHECK_FALSE(json_object_value(metadata, "properties").empty()); @@ -1070,9 +1071,15 @@ HRR_TEST_CASE(Unit_HRR_MetadataManifest) { CHECK(compute_mode >= 0); CHECK_FALSE(compute_capability.empty()); CHECK(std::count(compute_capability.begin(), compute_capability.end(), '.') == 1); - CHECK_FALSE(pci.empty()); - CHECK(std::count(pci.begin(), pci.end(), ':') == 2); - CHECK_FALSE(uuid.empty()); + // The keys stay, blank: they would identify the individual GPU and replay + // needs none of them. + for (const char* key : {"\"pci\": \"\"", "\"uuid\": \"\"", "\"luid\": \"\""}) { + INFO("Blank field: " << key); + CHECK(metadata.find(key) != std::string::npos); + } + CHECK(pci.empty()); + CHECK(uuid.empty()); + CHECK(luid.empty()); } HRR_TEST_CASE(Unit_HRR_StreamAdvancedRoundtrip) { From 6cb83d89e0d007b6a6c34a5d67c2f315d2c34f50 Mon Sep 17 00:00:00 2001 From: andraghetti Date: Sun, 18 Oct 2026 15:22:30 +0200 Subject: [PATCH 2/8] fix(hrr): verify allocated build-info sections and only registered fat binaries inspect_archive.py verify skipped .comment and .debug* sections carrying SHF_ALLOC. Capture leaves exactly those sections unchanged, so a build path in one of them was reported as clean. The verifier now checks them whatever their flags; the allocation guard stays in the capture-side scrubber. blobs/ also holds host buffer contents, and a buffer's first bytes can look like a fat binary: a 4-byte "CCOB" buffer counted as an unchecked compressed fat binary, and a buffer holding an offload bundle counted as a code object. The verifier now takes as fat binaries only the blobs that __hipRegisterFatBinary records in events.bin name by hash. Jira: AIRUNTIME-2881 --- .../scripts/inspect_archive.py | 66 ++++++++++++++----- .../tests/test_inspect_archive.py | 63 ++++++++++++++++-- 2 files changed, 106 insertions(+), 23 deletions(-) diff --git a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py index 541b1574a492..ef04d1c1a6d8 100755 --- a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py +++ b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py @@ -79,6 +79,11 @@ # See projects/clr/hipamd/src/hrr/hip_capture_scrub.h. BUNDLE_MAGIC = b"__CLANG_OFFLOAD_BUNDLE__" COMPRESSED_BUNDLE_MAGIC = b"CCOB" +# `blobs/` holds buffer contents as well as fat binaries, so the fat binaries +# are the blobs a `__hipRegisterFatBinary` record names: its hrr_api_id_t, and +# where the record keeps the hash after the event header and `uint64 ret`. +REGISTER_FAT_BINARY = 550 +FAT_BINARY_HASH_OFFSET = EVENT_HEADER_BYTES + 8 HARDWARE_ID_FIELDS = ("pci", "uuid", "luid") @@ -181,7 +186,11 @@ def _dir_stats(path: Path) -> tuple[int, int]: def _elf_has_build_info(data: bytes) -> bool: - """Does this ELF64 image hold a .comment or .debug* section that is not all zero?""" + """Does this ELF64 image hold a .comment or .debug* section that is not all zero? + + Allocated sections count too: capture leaves them as they are, so they are + exactly the ones that can still carry a build path. + """ if len(data) < 64 or data[:4] != b"\x7fELF" or data[4] != 2 or data[5] != 1: return False shoff = struct.unpack_from(" bool: return False names = struct.unpack_from("= 0 else len(data)] if label != b".comment" and not label.startswith(b".debug"): continue - if kind == 8 or flags & 0x2: # SHT_NOBITS, SHF_ALLOC + if kind == 8: # SHT_NOBITS: nothing in the file continue if data[offset : offset + size].strip(b"\0"): return True @@ -221,30 +230,55 @@ def _code_object_has_build_info(data: bytes) -> bool: return False +def _fat_binary_blobs(pid_dir: Path) -> list[Path]: + """The blob files that `__hipRegisterFatBinary` records in `events.bin` name. + + The first bytes of a blob cannot say this: a host buffer may well open with + "CCOB" or hold an offload bundle. Records are walked by their declared + length, as the reader does, up to the first one that is torn. + """ + events = pid_dir / "events.bin" + hashes: set[str] = set() + try: + size = events.stat().st_size + with events.open("rb") as handle: + head = handle.read(HEADER_BYTES) + if len(head) < HEADER_BYTES or struct.unpack_from(" min(size - pos, MAX_RECORD_BYTES): + break + if event_type == REGISTER_FAT_BINARY and length >= FAT_BINARY_HASH_OFFSET + 16: + record += handle.read(FAT_BINARY_HASH_OFFSET + 16 - EVENT_HEADER_BYTES) + lo, hi = struct.unpack_from(" None: """Count the stored code objects that still carry build information. Every file under `code_objects/` is one; under `blobs/` only the fat - binaries are, and the first bytes tell them from buffer contents without - reading the buffers. + binaries are, and the events say which those are. """ - candidates = sorted((pid_dir / "code_objects").glob("*")) + sorted( - (pid_dir / "blobs").rglob("*") - ) + candidates = sorted((pid_dir / "code_objects").glob("*")) + _fat_binary_blobs(pid_dir) for path in candidates: try: if not path.is_file(): continue - with path.open("rb") as handle: - head = handle.read(len(BUNDLE_MAGIC)) - if head.startswith(COMPRESSED_BUNDLE_MAGIC): - proc.unchecked_fat_binaries += 1 - continue - if path.parent.name != "code_objects" and head != BUNDLE_MAGIC: - continue - data = head + handle.read() + data = path.read_bytes() except OSError: continue + if data.startswith(COMPRESSED_BUNDLE_MAGIC): + proc.unchecked_fat_binaries += 1 + continue if _code_object_has_build_info(data): proc.build_info_files += 1 diff --git a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py index 05f7bbbe7ec7..ebf6061b3569 100644 --- a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py +++ b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py @@ -562,8 +562,11 @@ def test_json_output_is_machine_readable(tmp_path, capsys): assert payload["processes"][0]["pid"] == 99 -def make_elf(sections: dict[str, bytes]) -> bytes: - """A minimal ELF64 little-endian image holding `sections`, then .shstrtab.""" +def make_elf(sections: dict[str, bytes], allocated: tuple[str, ...] = ()) -> bytes: + """A minimal ELF64 little-endian image holding `sections`, then .shstrtab. + + Sections named in `allocated` carry SHF_ALLOC. + """ names = b"\0" offsets = {} for name in [*sections, ".shstrtab"]: @@ -573,7 +576,10 @@ def make_elf(sections: dict[str, bytes]) -> bytes: headers = [struct.pack(" bytes: return table + elf +def register_fat_binary(pid_dir: Path, lo: int, hi: int, data: bytes) -> None: + """Store `data` as the capture writer stores a fat binary, and append the + `__hipRegisterFatBinary` record that names it to `events.bin`. + """ + name = f"{lo:016x}{hi:016x}" + (pid_dir / "blobs" / name[:2]).mkdir(parents=True, exist_ok=True) + (pid_dir / "blobs" / name[:2] / f"{name}.blob").write_bytes(data) + record = struct.pack(" Date: Mon, 19 Oct 2026 06:16:37 +0200 Subject: [PATCH 3/8] fix(hrr): retype scrubbed code object sections SHT_NOBITS; bound the test's name read Zeroed .comment and .debug* sections are now SHT_NOBITS, so DWARF readers on a replay see no debug information instead of malformed zeros. They are not SHF_ALLOC and keep their offset, so the loader and the image size are unchanged. The test's ELF walker reads a section name only inside the string table. SKILL.md and README say what is removed: the debug information and the compiler build; the symbol table is not touched. --- .../clr/hipamd/src/hrr/hip_capture_scrub.cpp | 5 +- .../clr/hipamd/src/hrr/hip_capture_scrub.h | 3 +- projects/hrr/README.md | 4 +- projects/hrr/skills/enable-recording/SKILL.md | 2 +- .../integration/hrr_fatbin_header_test.cc | 82 +++++++++++++------ 5 files changed, 68 insertions(+), 28 deletions(-) diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp index 0996001e07f3..346a7a71e1af 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp @@ -71,7 +71,7 @@ size_t scrub_elf(uint8_t* image, size_t size) { size_t zeroed = 0; for (uint16_t i = 0; i < shnum; ++i) { - const uint8_t* sh = image + shoff + size_t{i} * kSectionHeaderSize; + uint8_t* sh = image + shoff + size_t{i} * kSectionHeaderSize; const uint32_t name = read(sh, 0); const uint32_t type = read(sh, 4); const uint64_t flags = read(sh, 8); @@ -82,6 +82,9 @@ size_t scrub_elf(uint8_t* image, size_t size) { if (!is_scrubbed_section(names + name, name_len)) continue; if (sec_size == 0 || !in_range(offset, sec_size, size)) continue; std::memset(image + offset, 0, sec_size); + // SHT_NOBITS makes DWARF readers see the section as absent rather than as + // malformed zeros. It is not SHF_ALLOC, so the loader maps nothing from it. + std::memcpy(sh + 4, &kShtNobits, sizeof(kShtNobits)); ++zeroed; } return zeroed; diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.h b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h index 84baf850c942..4d1fbdd3f719 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_scrub.h +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h @@ -15,7 +15,8 @@ namespace scrub { // object about to be stored: an ELF64 image, or every entry of an uncompressed // clang offload bundle. Those sections carry the compiler build, absolute build // paths and source file names; the GPU never loads them, so replay is unchanged. -// Sizes and offsets stay as they are, so a bundle's entry table stays valid. +// Each zeroed section is retyped SHT_NOBITS, so debuggers and profilers see it as +// absent. Sizes and offsets stay as they are, so a bundle's entry table stays valid. // A compressed bundle, or anything that does not parse, is left as it is. // Returns how many sections were zeroed. size_t scrub_code_object(uint8_t* image, size_t size); diff --git a/projects/hrr/README.md b/projects/hrr/README.md index ff3ea0edafd4..adfa6faacbd8 100644 --- a/projects/hrr/README.md +++ b/projects/hrr/README.md @@ -237,7 +237,9 @@ capture.hrr/ records the hash so kernels sharing a name (Triton emits many `triton_`) still resolve to the code object they came from. Their `.comment` and `.debug*` sections, which hold the compiler build and build paths, are - zeroed before they are stored, as are those of uncompressed fat binaries + 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. Symbol names stay - **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 diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index 808554476e77..4b230ae98280 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -130,7 +130,7 @@ What the archive holds, per process, and it is worth reading out to whoever has - **`events.bin`**: every successful HIP call the process made, with its arguments: sizes, device addresses, stream handles, kernel names. - **`blobs/`**: the host buffers those calls carried, byte for byte and in the clear. For a serving stack that means prompts and generated text; for training, the data and the weights that were copied to the device. -- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. Their `.comment` and `.debug*` sections are zeroed, here and in the fat binaries under `blobs/`, so the compiler build, source file names and build paths are gone, except in compressed (`CCOB`) fat binaries, which are stored as the application had them. Kernel names stay. +- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. Their `.comment` and `.debug*` sections are zeroed and marked empty, here and in the fat binaries under `blobs/`, which removes the compiler build and the build paths and source file names that debug information carries, except in compressed (`CCOB`) fat binaries, which are stored as the application had them. The symbol table is not touched, so kernel names stay. - **`manifest.json`**: process and parent process ids, event and blob counts, the clean-shutdown flag, the HIP and comgr versions that captured, and the device's name, architecture, memory size and clocks. Its PCI location, UUID and LUID are written blank. - **`regions/*.hrrr`**: the framework allocator's own layout, when a producer ran alongside capture. diff --git a/projects/hrr/tests/integration/hrr_fatbin_header_test.cc b/projects/hrr/tests/integration/hrr_fatbin_header_test.cc index 52fea190b970..e28d852f240a 100644 --- a/projects/hrr/tests/integration/hrr_fatbin_header_test.cc +++ b/projects/hrr/tests/integration/hrr_fatbin_header_test.cc @@ -428,43 +428,44 @@ TEST_CASE("Unit_HRR_CodeObjectScrub_Direct", "[.][hrr-direct]") { namespace { struct ScrubCount { size_t elves = 0; - size_t sections = 0; // .comment and .debug* sections with contents + size_t sections = 0; // non-empty .comment and .debug* sections size_t left = 0; // of those, still holding a non-zero byte + size_t typed = 0; // of those, not SHT_NOBITS, so DWARF readers parse them }; // Walks the section headers of an ELF64 image independently of capture's own -// parser and counts the sections capture has to zero. -void count_scrub_sections(const uint8_t* p, size_t size, ScrubCount& c) { +// parser and counts the sections capture has to scrub. False if the image does +// not parse, including a section name that does not end inside the string table. +bool count_scrub_sections(const uint8_t* p, size_t size, ScrubCount& c) { auto u16 = [&](size_t at) { uint16_t v; std::memcpy(&v, p + at, 2); return v; }; auto u32 = [&](size_t at) { uint32_t v; std::memcpy(&v, p + at, 4); return v; }; auto u64 = [&](size_t at) { uint64_t v; std::memcpy(&v, p + at, 8); return v; }; - REQUIRE(size >= 64); - REQUIRE(std::memcmp(p, "\x7f" "ELF", 4) == 0); - ++c.elves; + if (size < 64 || std::memcmp(p, "\x7f" "ELF", 4) != 0 || u16(0x3a) != 64) return false; const uint64_t shoff = u64(0x28); const uint16_t shnum = u16(0x3c); const uint16_t shstrndx = u16(0x3e); - REQUIRE(u16(0x3a) == 64); - REQUIRE(shoff + uint64_t{shnum} * 64 <= size); - REQUIRE(shstrndx < shnum); + if (shoff > size || uint64_t{shnum} * 64 > size - shoff || shstrndx >= shnum) return false; const uint64_t names = u64(shoff + shstrndx * 64 + 24); + const uint64_t names_size = u64(shoff + shstrndx * 64 + 32); + if (names > size || names_size > size - names) return false; + ++c.elves; for (uint16_t i = 0; i < shnum; ++i) { const size_t sh = shoff + size_t{i} * 64; - REQUIRE(names + u32(sh) < size); - const std::string name(reinterpret_cast(p + names + u32(sh))); + const uint32_t at = u32(sh); + if (at >= names_size) return false; + const char* s = reinterpret_cast(p + names + at); + const size_t name_len = strnlen(s, names_size - at); + if (name_len == names_size - at) return false; + const std::string name(s, name_len); if (name != ".comment" && name.rfind(".debug", 0) != 0) continue; const uint64_t off = u64(sh + 24), len = u64(sh + 32); - if (u32(sh + 4) == 8 /*SHT_NOBITS*/ || len == 0) continue; - REQUIRE(off + len <= size); - INFO("Section " << name); + if (len == 0) continue; + if (off > size || len > size - off) return false; ++c.sections; - for (uint64_t b = 0; b < len; ++b) { - if (p[off + b] != 0) { - ++c.left; - break; - } - } + if (u32(sh + 4) != 8 /*SHT_NOBITS*/) ++c.typed; + if (std::any_of(p + off, p + off + len, [](uint8_t b) { return b != 0; })) ++c.left; } + return true; } bool contains(const std::vector& bytes, const std::string& needle) { @@ -488,7 +489,7 @@ HRR_TEST_CASE(Unit_HRR_CodeObjectScrub) { const std::string text = read_text_file(entry.path()); const std::vector image(text.begin(), text.end()); ++module_files; - count_scrub_sections(image.data(), image.size(), module_count); + REQUIRE(count_scrub_sections(image.data(), image.size(), module_count)); CHECK_FALSE(contains(image, "clang version")); CHECK_FALSE(contains(image, "hrr_fatbin_probe.hip")); } @@ -511,21 +512,54 @@ HRR_TEST_CASE(Unit_HRR_CodeObjectScrub) { const uint64_t off = read_u64(blob, pos), len = read_u64(blob, pos + 8); pos += kEntryHeaderSize + read_u64(blob, pos + 16); if (len >= 4 && std::memcmp(blob.data() + off, "\x7f" "ELF", 4) == 0) - count_scrub_sections(blob.data() + off, len, bundle_count); + REQUIRE(count_scrub_sections(blob.data() + off, len, bundle_count)); } CHECK_FALSE(contains(blob, "clang version")); CHECK_FALSE(contains(blob, "hrr_fatbin_probe.hip")); } INFO("Module code objects: " << module_count.elves << " ELF, " << module_count.sections - << " sections, " << module_count.left << " not zeroed"); + << " sections, " << module_count.left << " not zeroed, " + << module_count.typed << " not SHT_NOBITS"); INFO("Fat binary code objects: " << bundle_count.elves << " ELF, " << bundle_count.sections - << " sections, " << bundle_count.left << " not zeroed"); + << " sections, " << bundle_count.left << " not zeroed, " + << bundle_count.typed << " not SHT_NOBITS"); // The probe was built with -g, so both copies had something to zero. CHECK(module_count.sections >= 2); CHECK(bundle_count.sections >= 2); CHECK(module_count.left == 0); CHECK(bundle_count.left == 0); + // Retyped, so a debugger on the replay sees no debug information rather than zeros. + CHECK(module_count.typed == 0); + CHECK(bundle_count.typed == 0); +} + +// The walker reads a section name only inside its string table, even when the +// bytes after the table would complete a name it looks for. +HRR_TEST_CASE(Unit_HRR_CodeObjectScrubWalker) { + std::vector elf(80 + 2 * 64, 0); + auto put = [&](size_t at, uint64_t v, size_t n) { std::memcpy(elf.data() + at, &v, n); }; + std::memcpy(elf.data(), "\x7f" "ELF", 4); + elf[4] = 2; // ELFCLASS64 + elf[5] = 1; // ELFDATA2LSB + put(0x28, 80, 8); // e_shoff + put(0x3a, 64, 2); // e_shentsize + put(0x3c, 2, 2); // e_shnum + put(0x3e, 1, 2); // e_shstrndx + std::memcpy(elf.data() + 64, "\0.comment", 10); + // Section 1 is the string table, named at offset 1, and declares only "\0.comm". + const size_t sh = 80 + 64; + put(sh, 1, 4); + put(sh + 4, 3 /*SHT_STRTAB*/, 4); + put(sh + 24, 64, 8); + put(sh + 32, 6, 8); + ScrubCount c; + CHECK_FALSE(count_scrub_sections(elf.data(), elf.size(), c)); + // Declared up to its terminator, the same table parses. + put(sh + 32, 10, 8); + c = ScrubCount{}; + CHECK(count_scrub_sections(elf.data(), elf.size(), c)); + CHECK(c.sections == 1); } #endif // HRR_TEST_EXE From b90e2f0fa674ec3c044100d1438df73f21571872 Mon Sep 17 00:00:00 2001 From: andraghetti Date: Mon, 19 Oct 2026 09:38:02 +0200 Subject: [PATCH 4/8] fix(hrr): skip unterminated section names; pick the fat binary id by archive version The scrubber matched a section name that has no terminator inside the string table, and the verifier read names past the table's end. Both now skip such a name. hrr-unit-tests compiles the scrubber on its own and feeds it malformed images. verify looked for __hipRegisterFatBinary under id 550 only, which is v9's. It now uses 549 for a v8 archive and reports fat binaries as not scanned for any other version. --- .../clr/hipamd/src/hrr/hip_capture_scrub.cpp | 1 + projects/hrr/skills/enable-recording/SKILL.md | 2 +- .../scripts/inspect_archive.py | 34 +++- .../tests/test_inspect_archive.py | 45 ++++- projects/hrr/tests/unit/CMakeLists.txt | 9 + projects/hrr/tests/unit/hrr_scrub_test.cc | 178 ++++++++++++++++++ 6 files changed, 256 insertions(+), 13 deletions(-) create mode 100644 projects/hrr/tests/unit/hrr_scrub_test.cc diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp index 346a7a71e1af..a774b55e76c2 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.cpp @@ -79,6 +79,7 @@ size_t scrub_elf(uint8_t* image, size_t size) { const uint64_t sec_size = read(sh, 32); if (name >= names_size || type == kShtNobits || (flags & kShfAlloc) != 0) continue; const size_t name_len = strnlen(names + name, names_size - name); + if (name_len == names_size - name) continue; // no terminator inside the table if (!is_scrubbed_section(names + name, name_len)) continue; if (sec_size == 0 || !in_range(offset, sec_size, size)) continue; std::memset(image + offset, 0, sec_size); diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index 4b230ae98280..6d396743f6f4 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -136,7 +136,7 @@ What the archive holds, per process, and it is worth reading out to whoever has Apart from the few exceptions in step 2, failed HIP calls are not recorded, and nothing that never crossed a HIP call is. The archive is exactly as sensitive as the data the workload handled, so where it may go has to be agreed before it is uploaded. -`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`, `build_info_files`, and `unchecked_fat_binaries` for the compressed ones it cannot read); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. +`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`, `build_info_files`, `unchecked_fat_binaries` for the compressed ones it cannot read, and `fat_binaries_scanned`, false when the archive's format version is not v8 or v9 and it cannot tell which blobs are fat binaries); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. If the content is the obstacle, three things exist and none of them is scrubbing. Re-capture the same failure driven by synthetic input, if it still reproduces. Replay the archive yourself and send only the finding, which is a few lines and carries no buffers. Or agree the handling of a production-data transfer before uploading, because that is what this is. diff --git a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py index ef04d1c1a6d8..f26b169ace85 100755 --- a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py +++ b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py @@ -82,7 +82,10 @@ # `blobs/` holds buffer contents as well as fat binaries, so the fat binaries # are the blobs a `__hipRegisterFatBinary` record names: its hrr_api_id_t, and # where the record keeps the hash after the event header and `uint64 ret`. -REGISTER_FAT_BINARY = 550 +# The id is per archive format version (hrr_api_args.h): v8 put two compiler +# APIs ahead of it and v9 one more. A v7 archive may hold either numbering, so +# that and any other version leaves the fat binaries unscanned. +REGISTER_FAT_BINARY = {8: 549, 9: 550} FAT_BINARY_HASH_OFFSET = EVENT_HEADER_BYTES + 8 HARDWARE_ID_FIELDS = ("pci", "uuid", "luid") @@ -111,6 +114,8 @@ class ProcessArchive: build_info_files: int = 0 # Compressed fat binaries, which cannot be checked without decompressing. unchecked_fat_binaries: int = 0 + # False when events.bin is in a format version whose fat binary id is unknown. + fat_binaries_scanned: bool = True @property def recorded(self) -> bool: @@ -197,11 +202,15 @@ def _elf_has_build_info(data: bytes) -> bool: shentsize, shnum, shstrndx = struct.unpack_from("= shnum or shoff + shnum * 64 > len(data): return False - names = struct.unpack_from(" len(data): + return False for i in range(shnum): name, kind, _, _, offset, size = struct.unpack_from("= 0 else len(data)] + end = data.find(b"\0", names + name, names + names_size) + if name >= names_size or end < 0: + continue # the name does not end inside the string table + label = data[names + name : end] if label != b".comment" and not label.startswith(b".debug"): continue if kind == 8: # SHT_NOBITS: nothing in the file @@ -230,8 +239,9 @@ def _code_object_has_build_info(data: bytes) -> bool: return False -def _fat_binary_blobs(pid_dir: Path) -> list[Path]: - """The blob files that `__hipRegisterFatBinary` records in `events.bin` name. +def _fat_binary_blobs(pid_dir: Path) -> list[Path] | None: + """The blob files that `__hipRegisterFatBinary` records in `events.bin` name, + or None when the format version does not say which id those records have. The first bytes of a blob cannot say this: a host buffer may well open with "CCOB" or hold an offload bundle. Records are walked by their declared @@ -245,13 +255,16 @@ def _fat_binary_blobs(pid_dir: Path) -> list[Path]: head = handle.read(HEADER_BYTES) if len(head) < HEADER_BYTES or struct.unpack_from(" min(size - pos, MAX_RECORD_BYTES): break - if event_type == REGISTER_FAT_BINARY and length >= FAT_BINARY_HASH_OFFSET + 16: + if event_type == register_fat_binary and length >= FAT_BINARY_HASH_OFFSET + 16: record += handle.read(FAT_BINARY_HASH_OFFSET + 16 - EVENT_HEADER_BYTES) lo, hi = struct.unpack_from(" None: Every file under `code_objects/` is one; under `blobs/` only the fat binaries are, and the events say which those are. """ - candidates = sorted((pid_dir / "code_objects").glob("*")) + _fat_binary_blobs(pid_dir) + fat_binaries = _fat_binary_blobs(pid_dir) + proc.fat_binaries_scanned = fat_binaries is not None + candidates = sorted((pid_dir / "code_objects").glob("*")) + (fat_binaries or []) for path in candidates: try: if not path.is_file(): @@ -831,6 +846,8 @@ def render(report: ArchiveReport) -> str: build_info = f"build info in {proc.build_info_files} code objects" if proc.unchecked_fat_binaries: build_info += f", {proc.unchecked_fat_binaries} compressed fat binaries not checked" + if not proc.fat_binaries_scanned: + build_info += f", fat binaries not checked (archive format v{proc.format_version})" lines.append(f" hardware identifiers: {hardware} {build_info}") incomplete = [p for p in report.processes if p.complete is False] @@ -899,6 +916,7 @@ def to_dict(report: ArchiveReport) -> dict: "hardware_ids": p.hardware_ids, "build_info_files": p.build_info_files, "unchecked_fat_binaries": p.unchecked_fat_binaries, + "fat_binaries_scanned": p.fat_binaries_scanned, } for p in report.processes ], diff --git a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py index ebf6061b3569..123093a607de 100644 --- a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py +++ b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py @@ -599,14 +599,16 @@ def make_bundle(elf: bytes) -> bytes: return table + elf -def register_fat_binary(pid_dir: Path, lo: int, hi: int, data: bytes) -> None: +def register_fat_binary( + pid_dir: Path, lo: int, hi: int, data: bytes, event_id: int = inspect_archive.REGISTER_FAT_BINARY[9] +) -> None: """Store `data` as the capture writer stores a fat binary, and append the `__hipRegisterFatBinary` record that names it to `events.bin`. """ name = f"{lo:016x}{hi:016x}" (pid_dir / "blobs" / name[:2]).mkdir(parents=True, exist_ok=True) (pid_dir / "blobs" / name[:2] / f"{name}.blob").write_bytes(data) - record = struct.pack(", +# so it is compiled in here and fed malformed images on its own. +set(HRR_CAPTURE_SRC ${CMAKE_CURRENT_SOURCE_DIR}/../../../clr/hipamd/src/hrr) +if(EXISTS ${HRR_CAPTURE_SRC}/hip_capture_scrub.cpp) + target_sources(hrr-unit-tests PRIVATE + hrr_scrub_test.cc ${HRR_CAPTURE_SRC}/hip_capture_scrub.cpp) + target_include_directories(hrr-unit-tests PRIVATE ${HRR_CAPTURE_SRC}) +endif() + # ../include makes the local harness headers (hrr_test_common.hh, # hrr_test_process.hh) available. hrr_reader.h, hip_playback.h, and # hrr/hrr_api_args.h come transitively from hrr::reader. diff --git a/projects/hrr/tests/unit/hrr_scrub_test.cc b/projects/hrr/tests/unit/hrr_scrub_test.cc new file mode 100644 index 000000000000..fdef7fec66d1 --- /dev/null +++ b/projects/hrr/tests/unit/hrr_scrub_test.cc @@ -0,0 +1,178 @@ +/* + * Copyright (c) Advanced Micro Devices, Inc., or its affiliates. + * + * SPDX-License-Identifier: MIT + */ + +/** + * @addtogroup HRR HRR Code Object Scrub + * @{ + * @ingroup HRRTest + * GPU-free unit tests for capture's code object scrubber (CLR's + * hip_capture_scrub.cpp, compiled in on its own). It parses the application's + * memory, so each case feeds it a malformed image and checks that it writes + * nothing outside a section it has fully validated. + */ + +#include "hrr_test_common.hh" +#include "hip_capture_scrub.h" + +#include +#include +#include +#include + +using hrr_cap::scrub::scrub_code_object; + +namespace { + +constexpr uint32_t kProgbits = 1, kStrtab = 3, kNobits = 8; + +template +void put(std::vector& v, size_t at, T x) { + std::memcpy(v.data() + at, &x, sizeof(x)); +} +template +T get(const std::vector& v, size_t at) { + T x; + std::memcpy(&x, v.data() + at, sizeof(x)); + return x; +} + +struct Section { + std::string name, data; +}; + +// ELF64 LSB: header, section contents, .shstrtab, then the section headers +// (null, `sections` in order, .shstrtab last). The last section's name is the +// last string in the table. +std::vector make_elf(const std::vector
& sections) { + std::string names = std::string(1, '\0') + ".shstrtab" + '\0'; + struct Header { uint32_t name, type; uint64_t offset, size; }; + std::vector
headers{{0, 0, 0, 0}}; + std::vector image(64, 0); + for (const auto& s : sections) { + headers.push_back({uint32_t(names.size()), kProgbits, image.size(), s.data.size()}); + names += s.name + '\0'; + image.insert(image.end(), s.data.begin(), s.data.end()); + } + headers.push_back({1, kStrtab, image.size(), names.size()}); + image.insert(image.end(), names.begin(), names.end()); + const uint64_t shoff = image.size(); + image.resize(shoff + headers.size() * 64, 0); + for (size_t i = 0; i < headers.size(); ++i) { + put(image, shoff + i * 64, headers[i].name); + put(image, shoff + i * 64 + 4, headers[i].type); + put(image, shoff + i * 64 + 24, headers[i].offset); + put(image, shoff + i * 64 + 32, headers[i].size); + } + std::memcpy(image.data(), "\x7f" "ELF\x02\x01", 6); // ELFCLASS64, ELFDATA2LSB + put(image, 0x28, shoff); + put(image, 0x3a, uint16_t{64}); + put(image, 0x3c, uint16_t(headers.size())); + put(image, 0x3e, uint16_t(headers.size() - 1)); + return image; +} + +size_t header_of(const std::vector& elf, size_t i) { + return get(elf, 0x28) + i * 64; +} + +std::string contents(const std::vector& elf, size_t i) { + const size_t sh = header_of(elf, i); + const auto* p = elf.data() + get(elf, sh + 24); + return std::string(p, p + get(elf, sh + 32)); +} + +// An uncompressed offload bundle with one entry, `elf`, whose id is `id_size` long. +std::vector make_bundle(const std::vector& elf, uint64_t id_size) { + const std::string id = "hipv4-amdgcn-amd-amdhsa--gfx950"; + std::vector b(24 + 8 + 24, 0); + std::memcpy(b.data(), "__CLANG_OFFLOAD_BUNDLE__", 24); + put(b, 24, uint64_t{1}); + put(b, 32, uint64_t(b.size() + id.size())); + put(b, 40, uint64_t(elf.size())); + put(b, 48, id_size); + b.insert(b.end(), id.begin(), id.end()); + b.insert(b.end(), elf.begin(), elf.end()); + return b; +} + +const std::vector
kLeaky{ + {".text", "\x01\x02\x03\x04"}, {".debug_str", "/app/build/k.hip"}, {".comment", "clang version"}}; + +} // namespace + +HRR_TEST_CASE(Unit_HRR_Scrub_ZeroesAndRetypesBuildInfo) { + auto elf = make_elf(kLeaky); + REQUIRE(scrub_code_object(elf.data(), elf.size()) == 2); + CHECK(contents(elf, 1) == kLeaky[0].data); + CHECK(get(elf, header_of(elf, 1) + 4) == kProgbits); + for (size_t i : {size_t{2}, size_t{3}}) { + CHECK(contents(elf, i) == std::string(kLeaky[i - 1].data.size(), '\0')); + CHECK(get(elf, header_of(elf, i) + 4) == kNobits); + } + + auto bundle = make_bundle(make_elf(kLeaky), 31); + REQUIRE(scrub_code_object(bundle.data(), bundle.size()) == 2); +} + +// The table declares ".comment" without its terminator, so the name is +// malformed even though the byte after the table is a NUL. +HRR_TEST_CASE(Unit_HRR_Scrub_SkipsUnterminatedName) { + auto elf = make_elf(kLeaky); + const size_t strtab = header_of(elf, 4); + put(elf, strtab + 32, get(elf, strtab + 32) - 1); + CHECK(scrub_code_object(elf.data(), elf.size()) == 1); // .debug_str only + CHECK(contents(elf, 3) == kLeaky[2].data); + CHECK(get(elf, header_of(elf, 3) + 4) == kProgbits); + CHECK(get(elf, header_of(elf, 2) + 4) == kNobits); +} + +HRR_TEST_CASE(Unit_HRR_Scrub_RefusesOutOfBoundsTables) { + const auto good = make_elf(kLeaky); + auto unchanged = [](std::vector elf) { + const auto before = elf; + return scrub_code_object(elf.data(), elf.size()) == 0 && elf == before; + }; + + // .comment runs one byte past the end of the image: only .debug_str changes. + auto elf = good; + put(elf, header_of(elf, 3) + 24, uint64_t(elf.size() - kLeaky[2].data.size() + 1)); + auto expected = elf; + const size_t debug = header_of(expected, 2); + std::memset(expected.data() + get(expected, debug + 24), 0, kLeaky[1].data.size()); + put(expected, debug + 4, kNobits); + CHECK(scrub_code_object(elf.data(), elf.size()) == 1); + CHECK(elf == expected); + + elf = good; // the string table runs past the end + put(elf, header_of(elf, 4) + 32, uint64_t(elf.size())); + CHECK(unchanged(elf)); + elf = good; // its offset wraps around + put(elf, header_of(elf, 4) + 24, ~uint64_t{0}); + CHECK(unchanged(elf)); + + elf = good; // the section table's offset wraps around + put(elf, 0x28, ~uint64_t{0} - 63); + CHECK(unchanged(elf)); + elf = good; // the section table ends one byte past the image + put(elf, 0x28, get(elf, 0x28) + 1); + CHECK(unchanged(elf)); + + // A bundle entry's id length runs past the end, and wraps the cursor if added. + CHECK(unchanged(make_bundle(good, ~uint64_t{0} - 40))); + CHECK(unchanged(make_bundle(good, 1u << 20))); + + // Every truncation of the image, and of the bundle, is refused. + const auto bundle = make_bundle(good, 31); + for (const auto* whole : {&good, &bundle}) + for (size_t n = 0; n < whole->size(); ++n) { + std::vector cut(whole->begin(), whole->begin() + n); + CHECK(scrub_code_object(cut.data(), cut.size()) == 0); + } +} + +/** + * @} + */ From 33c8ea625badb51ad16e042413ebf0cf08a1e1e3 Mon Sep 17 00:00:00 2001 From: andraghetti Date: Mon, 19 Oct 2026 09:48:32 +0200 Subject: [PATCH 5/8] docs(hrr): finish the README sentence on symbol names --- projects/hrr/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hrr/README.md b/projects/hrr/README.md index adfa6faacbd8..f03d13387ec8 100644 --- a/projects/hrr/README.md +++ b/projects/hrr/README.md @@ -239,7 +239,7 @@ capture.hrr/ `.debug*` sections, which hold the compiler build and build paths, 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. Symbol names stay + information. The symbol table is not touched, so kernel names stay. - **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 From 4ef0b5a56d210d7f748e610a1ba2329f2d8c1a8e Mon Sep 17 00:00:00 2001 From: andraghetti Date: Mon, 19 Oct 2026 12:02:45 +0200 Subject: [PATCH 6/8] fix(hrr): blank the GPU's UUID, LUID and PCI ids in hipChooseDevice records CLR's hipChooseDevice matches on numeric limits only, so replay needs none of them; the manifest already writes them blank. --- .../hipamd/src/hrr/hip_capture_generated.cpp | 9 ++++++ .../tests/integration/hrr_roundtrip_test.cc | 29 +++++++++++++++++++ projects/hrr/tools/gen_hrr_api_args.py | 16 ++++++++-- 3 files changed, 51 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp b/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp index db275464fc0a..833dd7793264 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_generated.cpp @@ -46,6 +46,7 @@ #include "utils/debug.hpp" #include +#include #include #include #include @@ -207,6 +208,11 @@ static hipError_t capture_hipChooseDevice(int* device, const hipDeviceProp_tR060 a.prop = reinterpret_cast(prop); if (prop) { std::memcpy(a.prop_bytes, prop, sizeof(hipDeviceProp_t)); + std::memset(a.prop_bytes + offsetof(hipDeviceProp_t, uuid), 0, sizeof(prop->uuid)); + std::memset(a.prop_bytes + offsetof(hipDeviceProp_t, luid), 0, sizeof(prop->luid)); + std::memset(a.prop_bytes + offsetof(hipDeviceProp_t, pciDomainID), 0, sizeof(prop->pciDomainID)); + std::memset(a.prop_bytes + offsetof(hipDeviceProp_t, pciBusID), 0, sizeof(prop->pciBusID)); + std::memset(a.prop_bytes + offsetof(hipDeviceProp_t, pciDeviceID), 0, sizeof(prop->pciDeviceID)); a.prop_present = 1; } hrr_cap::writer::write_event_raw(HRR_API_HIPCHOOSEDEVICE, &a.hdr, sizeof(a)); @@ -224,6 +230,9 @@ static hipError_t capture_hipChooseDeviceR0000(int* device, const hipDeviceProp_ a.properties = reinterpret_cast(properties); if (properties) { std::memcpy(a.properties_bytes, properties, sizeof(hipDeviceProp_tR0000)); + std::memset(a.properties_bytes + offsetof(hipDeviceProp_tR0000, pciDomainID), 0, sizeof(properties->pciDomainID)); + std::memset(a.properties_bytes + offsetof(hipDeviceProp_tR0000, pciBusID), 0, sizeof(properties->pciBusID)); + std::memset(a.properties_bytes + offsetof(hipDeviceProp_tR0000, pciDeviceID), 0, sizeof(properties->pciDeviceID)); a.properties_present = 1; } hrr_cap::writer::write_event_raw(HRR_API_HIPCHOOSEDEVICER0000, &a.hdr, sizeof(a)); diff --git a/projects/hrr/tests/integration/hrr_roundtrip_test.cc b/projects/hrr/tests/integration/hrr_roundtrip_test.cc index 8942dcfb1bc6..9998acdefe5a 100644 --- a/projects/hrr/tests/integration/hrr_roundtrip_test.cc +++ b/projects/hrr/tests/integration/hrr_roundtrip_test.cc @@ -1080,6 +1080,35 @@ HRR_TEST_CASE(Unit_HRR_MetadataManifest) { CHECK(pci.empty()); CHECK(uuid.empty()); CHECK(luid.empty()); + + // The workload hands hipChooseDevice the properties of device 0: the same + // identifiers are blanked in the copy events.bin keeps, and the rest stays. + uint16_t version = 0; + FILE* f = hrr::open_record_stream((archive_path / "events.bin").string(), HRR_MAGIC, + HRR_VERSION, &version); + REQUIRE(f != nullptr); + std::vector raw; + int choose_records = 0; + while (hrr::read_raw_record(f, raw) == hrr::RecordStatus::Ok) { + const auto* hdr = reinterpret_cast(raw.data()); + if (hdr->event_type != HRR_API_HIPCHOOSEDEVICE) continue; + REQUIRE(raw.size() >= sizeof(hrr_args_hipChooseDevice)); + const auto* a = reinterpret_cast(raw.data()); + REQUIRE(a->prop_present == 1); + hipDeviceProp_t prop; + std::memcpy(&prop, a->prop_bytes, sizeof(prop)); + ++choose_records; + CHECK(std::string(prop.name).size() > 0); + CHECK(prop.multiProcessorCount > 0); + CHECK(std::all_of(std::begin(prop.uuid.bytes), std::end(prop.uuid.bytes), + [](char c) { return c == 0; })); + CHECK(std::all_of(std::begin(prop.luid), std::end(prop.luid), [](char c) { return c == 0; })); + CHECK(prop.pciDomainID == 0); + CHECK(prop.pciBusID == 0); + CHECK(prop.pciDeviceID == 0); + } + fclose(f); + CHECK(choose_records == 1); } HRR_TEST_CASE(Unit_HRR_StreamAdvancedRoundtrip) { diff --git a/projects/hrr/tools/gen_hrr_api_args.py b/projects/hrr/tools/gen_hrr_api_args.py index 2dbbd07189bc..d6b72d22e9e2 100644 --- a/projects/hrr/tools/gen_hrr_api_args.py +++ b/projects/hrr/tools/gen_hrr_api_args.py @@ -1160,6 +1160,9 @@ class Deref: # the bytes field zero and records only whether the pointer was non-null. # The field keeps its size, so archives from before still read. capability: bool = False + # Members of the pointee that name the individual machine and that the API + # never reads: capture zeroes them in its copy, so replay passes zeros. + blank_members: Tuple[str, ...] = () @property def is_array(self) -> bool: @@ -1261,12 +1264,15 @@ def total_bytes(self) -> int: ], # The device properties being matched against are the whole question these # two ask. 1472 bytes is a large event, but a program calls hipChooseDevice - # once at startup, not in a loop. + # once at startup, not in a loop. CLR matches on numeric limits only, so + # the GPU's UUID, LUID and PCI location are blanked, as in the manifest. "hipChooseDevice": [ - Deref("prop", "hipDeviceProp_t", 1472), + Deref("prop", "hipDeviceProp_t", 1472, + blank_members=("uuid", "luid", "pciDomainID", "pciBusID", "pciDeviceID")), ], "hipChooseDeviceR0000": [ - Deref("properties", "hipDeviceProp_tR0000", 792), + Deref("properties", "hipDeviceProp_tR0000", 792, + blank_members=("pciDomainID", "pciBusID", "pciDeviceID")), ], # The format decides the answer: the maximum width depends on the element # size the caller asked about. @@ -2380,6 +2386,7 @@ def generate_header(entries: List[ApiEntry]) -> str: #include "utils/debug.hpp" #include +#include #include #include #include @@ -2578,6 +2585,9 @@ def _fill_derefs(lines: List[str], entry: ApiEntry) -> None: else: lines.append(f" if ({d.param}) {{") lines.append(f" std::memcpy(a.{d.bytes_field}, {d.param}, sizeof({d.ctype}));") + for m in d.blank_members: + lines.append(f" std::memset(a.{d.bytes_field} + offsetof({d.ctype}, {m}), 0," + f" sizeof({d.param}->{m}));") lines.append(f" a.{d.present_field} = 1;") lines.append(f" }}") From 336c4ddb8770bea0f197902b1a0abe52197b38e2 Mon Sep 17 00:00:00 2001 From: andraghetti Date: Mon, 19 Oct 2026 12:02:45 +0200 Subject: [PATCH 7/8] fix(hrr): report absolute paths left in loaded code; say what capture keeps A device assert() keeps its source path in .rodata, which capture cannot zero. verify now counts such code objects (path_files) and calls the .comment/.debug* count debug_info_files. The docs and PRIVACY.md say what stays: loaded strings, compressed fat binaries, hipLinkAddData bitcode, and luid_device_node_mask. --- projects/clr/hipamd/src/hrr/hip_capture.cpp | 2 +- .../clr/hipamd/src/hrr/hip_capture_scrub.h | 1 + .../clr/hipamd/src/hrr/hip_capture_writer.cpp | 4 +- projects/hrr/PRIVACY.md | 14 +++-- projects/hrr/README.md | 11 ++-- projects/hrr/skills/enable-recording/SKILL.md | 10 +-- .../scripts/inspect_archive.py | 62 ++++++++++++------- .../tests/test_inspect_archive.py | 44 ++++++++++--- 8 files changed, 99 insertions(+), 49 deletions(-) diff --git a/projects/clr/hipamd/src/hrr/hip_capture.cpp b/projects/clr/hipamd/src/hrr/hip_capture.cpp index 2a88c4f4d7a0..ae91276c8eba 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture.cpp @@ -1444,7 +1444,7 @@ static size_t compute_bundle_size(const void* blob) { return static_cast(end); } -// Stores a fat binary with the build paths and compiler build of its code +// Stores a fat binary with the .comment and debug sections of its code // objects zeroed; the application's own copy is left alone. static hrr_cap::Hash128 write_fat_binary_blob(const void* blob, size_t blob_size) { std::vector scrubbed(static_cast(blob), diff --git a/projects/clr/hipamd/src/hrr/hip_capture_scrub.h b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h index 4d1fbdd3f719..e02026d17a20 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_scrub.h +++ b/projects/clr/hipamd/src/hrr/hip_capture_scrub.h @@ -15,6 +15,7 @@ namespace scrub { // object about to be stored: an ELF64 image, or every entry of an uncompressed // clang offload bundle. Those sections carry the compiler build, absolute build // paths and source file names; the GPU never loads them, so replay is unchanged. +// A string the GPU does load, such as the file name in a device assert(), stays. // Each zeroed section is retyped SHT_NOBITS, so debuggers and profilers see it as // absent. Sizes and offsets stay as they are, so a bundle's entry table stays valid. // A compressed bundle, or anything that does not parse, is left as it is. diff --git a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp index d72b20324a57..2dd1c54638ce 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp @@ -1630,8 +1630,8 @@ Hash128 write_code_object(const void* image, size_t image_size) { if (g_events_fd < 0 || g_events_finalized) return {}; // not open, or past the trailer } - // The stored bytes, and so the hash, are the image without its build paths - // and compiler build (see hip_capture_scrub.h). + // The stored bytes, and so the hash, are the image with its .comment and + // debug sections zeroed (see hip_capture_scrub.h). std::vector scrubbed(static_cast(image), static_cast(image) + image_size); scrub::scrub_code_object(scrubbed.data(), scrubbed.size()); diff --git a/projects/hrr/PRIVACY.md b/projects/hrr/PRIVACY.md index 7c9af5dbf2e5..6cb5e6e1949a 100644 --- a/projects/hrr/PRIVACY.md +++ b/projects/hrr/PRIVACY.md @@ -48,15 +48,16 @@ every capture is a contents capture. | Sequence of successful HIP calls, with sizes, flags, device addresses, handles, thread ids | `events.bin` | both | — kept, replay needs it | | Kernel names, mangled but read back with `c++filt`; they describe the model's architecture and precision | `events.bin` | both | — kept, triage names the failing kernel | | Event timestamps, from a monotonic clock (not the date or time of day) | `events.bin` | both | — kept | -| Device properties passed to `hipChooseDevice`, the whole `hipDeviceProp_t`: device and architecture names, PCI ids and, in the current layout, the UUID and LUID | `events.bin` | both | — kept, the call's answer depends on them | +| Device properties passed to `hipChooseDevice`, the whole `hipDeviceProp_t` with the device and architecture names; its UUID, LUID and PCI ids are written as zeros | `events.bin` | both | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): UUID, LUID and PCI ids blanked; the rest is kept, the call's answer depends on it | | Other structures passed into a HIP call, copied whole, such as copy, memset, graph node, allocation, pool and launch-attribute parameters. Padding, reserved and unused union bytes are copied as the caller left them | `events.bin` | both | — kept, replay needs it | | 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) | -| Linker input buffers passed to `hipLinkAddData`, every byte, with their names, including inputs the call rejected and inputs never linked or loaded | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | +| 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): stripped | +| 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 | | 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 | @@ -69,10 +70,11 @@ The per-process `manifest.json` also holds `pid`, `parent_pid`, `complete`, `eve | `total_global_mem`, `multi_processor_count`, `warp_size`, `max_threads_per_block` | yes | — | | `clock_rate_khz`, `memory_clock_rate_khz`, `memory_bus_width`, `l2_cache_size` | yes | — | | `integrated`, `managed_memory`, `memory_pools_supported`, `compute_mode` | yes | — | -| `uuid`, `pci`, `luid`, `luid_device_node_mask`: identify the individual GPU and its slot | no | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): blanked | +| `luid_device_node_mask`, a node mask rather than an identifier | yes | — | +| `uuid`, `pci`, `luid`: identify the individual GPU and its slot | written blank | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881) | -This table covers `manifest.json` only. A `hipChooseDevice` call can carry the same -identifiers in `events.bin`, as listed above. +This table covers `manifest.json` only. The copy of the device properties a +`hipChooseDevice` call keeps in `events.bin` has the same identifiers blanked, as listed above. ## Never recorded diff --git a/projects/hrr/README.md b/projects/hrr/README.md index c01ba28aa3d8..1e36041a5300 100644 --- a/projects/hrr/README.md +++ b/projects/hrr/README.md @@ -236,10 +236,13 @@ capture.hrr/ extracted by the runtime, for `hipModuleLoad` too); a launch records the hash so kernels sharing a name (Triton emits many `triton_`) still resolve to the code object they came from. Their `.comment` and - `.debug*` sections, which hold the compiler build and build paths, 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. The symbol table is not touched, so kernel names stay. + `.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. - **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 diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index 17b5b9cb6535..0721455020fb 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -124,19 +124,19 @@ An archive with several `pid-*` directories is one archive. Send the whole direc ## 4. Share it, knowing what is in it -**There is no scrubbing, anonymisation or randomisation option for the data today.** Say that plainly before anyone asks whether the archive can be sanitised, because the answer shapes what they do next and there is no flag that changes it. Capture does drop two things on its own, with no flag: the build paths and compiler build inside the code objects, and the GPU's hardware identifiers. +**There is no scrubbing, anonymisation or randomisation option for the data today.** Say that plainly before anyone asks whether the archive can be sanitised, because the answer shapes what they do next and there is no flag that changes it. Capture does drop two things on its own, with no flag: the `.comment` and debug sections of the code objects, which hold the compiler build and most build paths, and the GPU's hardware identifiers. What the archive holds, per process, and it is worth reading out to whoever has to approve the transfer: -- **`events.bin`**: every successful HIP call the process made, with its arguments: sizes, device addresses, stream handles, kernel names. -- **`blobs/`**: the host buffers those calls carried, byte for byte and in the clear. For a serving stack that means prompts and generated text; for training, the data and the weights that were copied to the device. -- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. Their `.comment` and `.debug*` sections are zeroed and marked empty, here and in the fat binaries under `blobs/`, which removes the compiler build and the build paths and source file names that debug information carries, except in compressed (`CCOB`) fat binaries, which are stored as the application had them. The symbol table is not touched, so kernel names stay. +- **`events.bin`**: every successful HIP call the process made, with its arguments: sizes, device addresses, stream handles, kernel names. The properties handed to `hipChooseDevice` are kept with their UUID, LUID and PCI location written as zeros. +- **`blobs/`**: the host buffers those calls carried, byte for byte and in the clear. For a serving stack that means prompts and generated text; for training, the data and the weights that were copied to the device. Linker inputs given to `hipLinkAddData` are stored as they were passed: SPIR-V or LLVM bitcode keeps its compiler identification and, built with `-g`, its source and build directories. +- **`code_objects/`**: the compiled GPU kernels the process loaded, which is the application's own binary code and a separate disclosure question from customer data. Their `.comment` and `.debug*` sections are zeroed and marked empty, here and in the fat binaries under `blobs/`, which removes the compiler build and the build paths and source file names that debug information carries, except in compressed (`CCOB`) fat binaries, which are stored as the application had them. What the GPU loads is not touched: kernel names stay, and so does any string the kernels keep, such as the source path a device `assert()` prints. - **`manifest.json`**: process and parent process ids, event and blob counts, the clean-shutdown flag, the HIP and comgr versions that captured, and the device's name, architecture, memory size and clocks. Its PCI location, UUID and LUID are written blank. - **`regions/*.hrrr`**: the framework allocator's own layout, when a producer ran alongside capture. Apart from the few exceptions in step 2, failed HIP calls are not recorded, and nothing that never crossed a HIP call is. The archive is exactly as sensitive as the data the workload handled, so where it may go has to be agreed before it is uploaded. [PRIVACY.md](../../PRIVACY.md) is the page to hand to whoever approves it. -`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`, `build_info_files`, `unchecked_fat_binaries` for the compressed ones it cannot read, and `fat_binaries_scanned`, false when the archive's format version is not v8 or v9 and it cannot tell which blobs are fat binaries); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. +`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`; `debug_info_files` for code objects still holding `.comment` or `.debug*` bytes; `path_files` for code objects whose loaded data holds an absolute path; `unchecked_fat_binaries` for the compressed ones it cannot read; `fat_binaries_scanned`, false when the archive's format version is not v8 or v9 and it cannot tell which blobs are fat binaries); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. If the content is the obstacle, three things exist and none of them is scrubbing. Re-capture the same failure driven by synthetic input, if it still reproduces. Replay the archive yourself and send only the finding, which is a few lines and carries no buffers. Or agree the handling of a production-data transfer before uploading, because that is what this is. diff --git a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py index f26b169ace85..36595310fe91 100755 --- a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py +++ b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py @@ -28,6 +28,7 @@ import argparse import json import os +import re import shutil import struct import subprocess @@ -88,6 +89,11 @@ REGISTER_FAT_BINARY = {8: 549, 9: 550} FAT_BINARY_HASH_OFFSET = EVENT_HEADER_BYTES + 8 HARDWARE_ID_FIELDS = ("pci", "uuid", "luid") +# Capture cannot zero what the GPU loads, so `verify` looks there for strings +# such as "/home/dev/src/model.hip": a NUL-terminated string that starts with +# "/" and has at least two components, in a loaded, non-executable section. +SHT_PROGBITS, SHF_ALLOC, SHF_EXECINSTR = 1, 0x2, 0x4 +ABSOLUTE_PATH = re.compile(rb"(?<=\0)/[\x21-\x7e]+/[\x21-\x7e]+(?=\0)") @dataclass @@ -111,7 +117,9 @@ class ProcessArchive: # Manifest fields that name the individual GPU and are not blank. hardware_ids: list[str] = field(default_factory=list) # Code objects and fat binaries still holding a .comment or .debug* section. - build_info_files: int = 0 + debug_info_files: int = 0 + # Code objects and fat binaries whose loaded data holds an absolute path. + path_files: int = 0 # Compressed fat binaries, which cannot be checked without decompressing. unchecked_fat_binaries: int = 0 # False when events.bin is in a format version whose fat binary id is unknown. @@ -190,23 +198,28 @@ def _dir_stats(path: Path) -> tuple[int, int]: return files, total -def _elf_has_build_info(data: bytes) -> bool: - """Does this ELF64 image hold a .comment or .debug* section that is not all zero? +def _elf_findings(data: bytes) -> tuple[bool, bool]: + """What an ELF64 image still carries: (a .comment or .debug* section that + is not all zero, an absolute path string in its loaded data). Allocated sections count too: capture leaves them as they are, so they are - exactly the ones that can still carry a build path. + exactly the ones that can still carry a build path, such as the source path + a device assert() prints, which sits in .rodata. """ + debug_info = path = False if len(data) < 64 or data[:4] != b"\x7fELF" or data[4] != 2 or data[5] != 1: - return False + return debug_info, path shoff = struct.unpack_from("= shnum or shoff + shnum * 64 > len(data): - return False + return debug_info, path names, names_size = struct.unpack_from(" len(data): - return False + return debug_info, path for i in range(shnum): - name, kind, _, _, offset, size = struct.unpack_from("= names_size or end < 0: continue # the name does not end inside the string table @@ -215,18 +228,18 @@ def _elf_has_build_info(data: bytes) -> bool: continue if kind == 8: # SHT_NOBITS: nothing in the file continue - if data[offset : offset + size].strip(b"\0"): - return True - return False + debug_info = debug_info or bool(data[offset : offset + size].strip(b"\0")) + return debug_info, path -def _code_object_has_build_info(data: bytes) -> bool: - """The same, for an ELF or for any entry of an uncompressed offload bundle.""" +def _code_object_findings(data: bytes) -> tuple[bool, bool]: + """The same, for an ELF or across the entries of an uncompressed offload bundle.""" if not data.startswith(BUNDLE_MAGIC): - return _elf_has_build_info(data) + return _elf_findings(data) + debug_info = path = False pos = len(BUNDLE_MAGIC) if len(data) < pos + 8: - return False + return debug_info, path (count,) = struct.unpack_from(" bool: break offset, size, id_size = struct.unpack_from(" list[Path] | None: @@ -294,8 +307,9 @@ def _scan_code_objects(pid_dir: Path, proc: ProcessArchive) -> None: if data.startswith(COMPRESSED_BUNDLE_MAGIC): proc.unchecked_fat_binaries += 1 continue - if _code_object_has_build_info(data): - proc.build_info_files += 1 + debug_info, path = _code_object_findings(data) + proc.debug_info_files += debug_info + proc.path_files += path def _archive_format_version(events: Path) -> int | None: @@ -843,7 +857,10 @@ def render(report: ArchiveReport) -> str: if proc.parent_pid: lines.append(f" parent pid: {proc.parent_pid}") hardware = ", ".join(proc.hardware_ids) if proc.hardware_ids else "blank" - build_info = f"build info in {proc.build_info_files} code objects" + build_info = ( + f".comment/.debug* in {proc.debug_info_files} code objects, " + f"absolute paths in {proc.path_files}" + ) if proc.unchecked_fat_binaries: build_info += f", {proc.unchecked_fat_binaries} compressed fat binaries not checked" if not proc.fat_binaries_scanned: @@ -914,7 +931,8 @@ def to_dict(report: ArchiveReport) -> dict: "code_object_files": p.code_object_files, "finalized": p.finalized, "hardware_ids": p.hardware_ids, - "build_info_files": p.build_info_files, + "debug_info_files": p.debug_info_files, + "path_files": p.path_files, "unchecked_fat_binaries": p.unchecked_fat_binaries, "fat_binaries_scanned": p.fat_binaries_scanned, } diff --git a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py index 123093a607de..2b2417d960db 100644 --- a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py +++ b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py @@ -645,19 +645,19 @@ def test_verify_reports_hardware_identifiers_and_build_info(tmp_path, capsys): inspect_archive.main(["--archive", str(tmp_path), "--no-playback", "--json"]) procs = {p["pid"]: p for p in json.loads(capsys.readouterr().out)["processes"]} assert procs[31]["hardware_ids"] == ["pci", "uuid"] - assert procs[31]["build_info_files"] == 2 + assert procs[31]["debug_info_files"] == 2 assert procs[31]["unchecked_fat_binaries"] == 1 assert procs[32]["hardware_ids"] == [] - assert procs[32]["build_info_files"] == 0 + assert procs[32]["debug_info_files"] == 0 assert procs[32]["unchecked_fat_binaries"] == 0 inspect_archive.main(["--archive", str(tmp_path), "--no-playback"]) text = capsys.readouterr().out assert ( - "hardware identifiers: pci, uuid build info in 2 code objects, " - "1 compressed fat binaries not checked" in text + "hardware identifiers: pci, uuid .comment/.debug* in 2 code objects, " + "absolute paths in 0, 1 compressed fat binaries not checked" in text ) - assert "hardware identifiers: blank build info in 0 code objects" in text + assert "hardware identifiers: blank .comment/.debug* in 0 code objects, absolute paths in 0" in text def test_an_allocated_comment_still_counts_as_build_info(tmp_path, capsys): @@ -670,7 +670,33 @@ def test_an_allocated_comment_still_counts_as_build_info(tmp_path, capsys): make_elf({".comment": b"/home/dev/build\0"}, allocated=(".comment",)) ) inspect_archive.main(["--archive", str(tmp_path), "--no-playback", "--json"]) - assert json.loads(capsys.readouterr().out)["processes"][0]["build_info_files"] == 1 + assert json.loads(capsys.readouterr().out)["processes"][0]["debug_info_files"] == 1 + + +def test_an_absolute_path_in_loaded_data_is_reported(tmp_path, capsys): + """A device assert() keeps its source file name in .rodata, which the GPU + loads, so capture leaves it and `verify` has to say it is there. Machine + code and sections the GPU never loads are not read as strings. + """ + leaky = make_process(tmp_path, 35) + (leaky / "code_objects").mkdir() + (leaky / "code_objects" / "a.hsaco").write_bytes( + make_elf({".rodata": b"p >= 0\0/home/dev/src/k.hip\0"}, allocated=(".rodata",)) + ) + clean = make_process(tmp_path, 36) + (clean / "code_objects").mkdir() + (clean / "code_objects" / "a.hsaco").write_bytes( + make_elf({".rodata": b"p >= 0\0k.hip\0//\0", ".note": b"\0/a/b\0"}, allocated=(".rodata",)) + ) + text = bytearray(make_elf({".text": b"\0/a/b\0"}, allocated=(".text",))) + struct.pack_into(" Date: Tue, 20 Oct 2026 17:53:30 +0200 Subject: [PATCH 8/8] fix(hrr): verify reads format v10 and paths with spaces or a drive letter; keep scrub cases listed without CLR inspect_archive.py had no __hipRegisterFatBinary id for format v10, so every archive develop writes skipped the fat binary scan. Map v10 to 550, and fail hrr-generator-rules (tools/tests/test_skill_api_ids.py) when HRR_VERSION in hrr_api_args.h moves past the table. The absolute path pattern now takes any byte from 0x20 up and a drive letter prefix, so paths with spaces, UTF-8 names and C:\ paths are reported. hrr_scrub_test.cc registers its three cases as skipped when the CLR sources are not checked out, so the case list still matches tests/config/unit.yaml. The case set is unchanged. --- projects/hrr/skills/enable-recording/SKILL.md | 2 +- .../scripts/inspect_archive.py | 18 ++++--- .../tests/test_inspect_archive.py | 36 ++++++++++--- projects/hrr/tests/CMakeLists.txt | 3 +- projects/hrr/tests/unit/CMakeLists.txt | 9 ++-- projects/hrr/tests/unit/hrr_scrub_test.cc | 11 ++++ .../hrr/tools/tests/test_skill_api_ids.py | 51 +++++++++++++++++++ 7 files changed, 112 insertions(+), 18 deletions(-) create mode 100644 projects/hrr/tools/tests/test_skill_api_ids.py diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index 0721455020fb..cf87598eca9d 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -136,7 +136,7 @@ What the archive holds, per process, and it is worth reading out to whoever has Apart from the few exceptions in step 2, failed HIP calls are not recorded, and nothing that never crossed a HIP call is. The archive is exactly as sensitive as the data the workload handled, so where it may go has to be agreed before it is uploaded. [PRIVACY.md](../../PRIVACY.md) is the page to hand to whoever approves it. -`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`; `debug_info_files` for code objects still holding `.comment` or `.debug*` bytes; `path_files` for code objects whose loaded data holds an absolute path; `unchecked_fat_binaries` for the compressed ones it cannot read; `fat_binaries_scanned`, false when the archive's format version is not v8 or v9 and it cannot tell which blobs are fat binaries); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. +`verify --json` gives each process's counts and sizes in a form you can attach to the request, and says whether any hardware identifier or code object build information is left (`hardware_ids`; `debug_info_files` for code objects still holding `.comment` or `.debug*` bytes; `path_files` for code objects whose loaded data holds an absolute path; `unchecked_fat_binaries` for the compressed ones it cannot read; `fat_binaries_scanned`, false when the archive's format version is not v8 to v10 and it cannot tell which blobs are fat binaries); the device details are in each `manifest.json`. For a multi-process archive, `hrr-playback --info` against the archive root prints a table covering every process at once. The cross-check in `verify` reports on one process, which is enough to judge the capture and not enough for a disclosure inventory. If the content is the obstacle, three things exist and none of them is scrubbing. Re-capture the same failure driven by synthetic input, if it still reproduces. Replay the archive yourself and send only the finding, which is a few lines and carries no buffers. Or agree the handling of a production-data transfer before uploading, because that is what this is. diff --git a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py index 36595310fe91..2b39f5ca00ed 100755 --- a/projects/hrr/skills/enable-recording/scripts/inspect_archive.py +++ b/projects/hrr/skills/enable-recording/scripts/inspect_archive.py @@ -84,16 +84,22 @@ # are the blobs a `__hipRegisterFatBinary` record names: its hrr_api_id_t, and # where the record keeps the hash after the event header and `uint64 ret`. # The id is per archive format version (hrr_api_args.h): v8 put two compiler -# APIs ahead of it and v9 one more. A v7 archive may hold either numbering, so -# that and any other version leaves the fat binaries unscanned. -REGISTER_FAT_BINARY = {8: 549, 9: 550} +# APIs ahead of it, v9 one more, and v10 left it where it was. A v7 archive may +# hold either numbering, so that and any other version leaves the fat binaries +# unscanned. tools/tests/test_skill_api_ids.py fails when HRR_VERSION moves +# past this table. +REGISTER_FAT_BINARY = {8: 549, 9: 550, 10: 550} FAT_BINARY_HASH_OFFSET = EVENT_HEADER_BYTES + 8 HARDWARE_ID_FIELDS = ("pci", "uuid", "luid") # Capture cannot zero what the GPU loads, so `verify` looks there for strings -# such as "/home/dev/src/model.hip": a NUL-terminated string that starts with -# "/" and has at least two components, in a loaded, non-executable section. +# such as "/home/dev/src/model.hip" or "C:\src\model.hip": a NUL-terminated +# string that starts with "/" or a drive letter and has at least two components, +# in a loaded, non-executable section. A component is any bytes from 0x20 up, +# so spaces and UTF-8 names count. SHT_PROGBITS, SHF_ALLOC, SHF_EXECINSTR = 1, 0x2, 0x4 -ABSOLUTE_PATH = re.compile(rb"(?<=\0)/[\x21-\x7e]+/[\x21-\x7e]+(?=\0)") +ABSOLUTE_PATH = re.compile( + rb"(?<=\0)(?:/[\x20-\xff]+/|[A-Za-z]:[\\/][\x20-\xff]+[\\/])[\x20-\xff]+(?=\0)" +) @dataclass diff --git a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py index 2b2417d960db..5fa2eae6838d 100644 --- a/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py +++ b/projects/hrr/skills/enable-recording/tests/test_inspect_archive.py @@ -686,7 +686,7 @@ def test_an_absolute_path_in_loaded_data_is_reported(tmp_path, capsys): clean = make_process(tmp_path, 36) (clean / "code_objects").mkdir() (clean / "code_objects" / "a.hsaco").write_bytes( - make_elf({".rodata": b"p >= 0\0k.hip\0//\0", ".note": b"\0/a/b\0"}, allocated=(".rodata",)) + make_elf({".rodata": b"p >= 0\0k.hip\0//\0C:\\k.hip\0", ".note": b"\0/a/b\0"}, allocated=(".rodata",)) ) text = bytearray(make_elf({".text": b"\0/a/b\0"}, allocated=(".text",))) struct.pack_into("= 0\0" + string + b"\0"}, allocated=(".rodata",)) + ) + inspect_archive.main(["--archive", str(tmp_path), "--no-playback", "--json"]) + assert json.loads(capsys.readouterr().out)["processes"][0]["path_files"] == 1 + + def test_only_blobs_registered_as_fat_binaries_are_checked(tmp_path, capsys): """`blobs/` also holds host buffers, whose bytes can be anything: a 4-byte buffer "CCOB", or an offload bundle the application keeps in memory. Only @@ -718,19 +740,19 @@ def test_only_blobs_registered_as_fat_binaries_are_checked(tmp_path, capsys): def test_fat_binaries_are_found_by_the_id_of_the_archive_version(tmp_path, capsys): - """`__hipRegisterFatBinary` was id 549 in format v8 and is 550 from v9. A v7 - archive may hold either, so `verify` says it did not look rather than - report no build information. + """`__hipRegisterFatBinary` was id 549 in format v8 and is 550 in v9 and + v10. A v7 archive may hold either, so `verify` says it did not look rather + than report no build information. """ leaky = make_bundle(make_elf({".text": b"\1" * 8, ".debug_str": b"/app/build\0"})) - cases = {40: (8, 549), 41: (8, 550), 42: (9, 550), 43: (7, 549)} + cases = {40: (8, 549), 41: (8, 550), 42: (9, 550), 43: (10, 550), 39: (7, 549)} for pid, (version, event_id) in cases.items(): register_fat_binary(make_process(tmp_path, pid, format_version=version), pid, 1, leaky, event_id) inspect_archive.main(["--archive", str(tmp_path), "--no-playback", "--json"]) procs = {p["pid"]: p for p in json.loads(capsys.readouterr().out)["processes"]} - assert [procs[pid]["debug_info_files"] for pid in cases] == [1, 0, 1, 0] - assert [procs[pid]["fat_binaries_scanned"] for pid in cases] == [True, True, True, False] + assert [procs[pid]["debug_info_files"] for pid in cases] == [1, 0, 1, 1, 0] + assert [procs[pid]["fat_binaries_scanned"] for pid in cases] == [True, True, True, True, False] inspect_archive.main(["--archive", str(tmp_path), "--no-playback"]) assert "fat binaries not checked (archive format v7)" in capsys.readouterr().out diff --git a/projects/hrr/tests/CMakeLists.txt b/projects/hrr/tests/CMakeLists.txt index 88009139715f..52cbe632ef3c 100644 --- a/projects/hrr/tests/CMakeLists.txt +++ b/projects/hrr/tests/CMakeLists.txt @@ -45,7 +45,8 @@ add_subdirectory(unit) # Also fails if a checked-in generated source or the API matrix header is # stale against the generator (test_generated_files_current.py), and if # PRIVACY.md no longer lists the manifest fields capture writes -# (test_privacy_doc.py). +# (test_privacy_doc.py), and if the enable-recording skill does not know the +# current archive format version (test_skill_api_ids.py). find_package(Python3 COMPONENTS Interpreter) if(Python3_Interpreter_FOUND) add_test(NAME hrr-generator-rules diff --git a/projects/hrr/tests/unit/CMakeLists.txt b/projects/hrr/tests/unit/CMakeLists.txt index 184ba263376a..ba40a956a81b 100644 --- a/projects/hrr/tests/unit/CMakeLists.txt +++ b/projects/hrr/tests/unit/CMakeLists.txt @@ -22,12 +22,15 @@ set_target_properties(hrr-unit-tests PROPERTIES CXX_STANDARD_REQUIRED ON) # Capture's code object scrubber lives in CLR and needs nothing but , -# so it is compiled in here and fed malformed images on its own. +# so it is compiled in here and fed malformed images on its own. A checkout of +# projects/hrr without the CLR sources still registers its cases, as skipped, +# so the binary keeps matching tests/config/unit.yaml. +target_sources(hrr-unit-tests PRIVATE hrr_scrub_test.cc) set(HRR_CAPTURE_SRC ${CMAKE_CURRENT_SOURCE_DIR}/../../../clr/hipamd/src/hrr) if(EXISTS ${HRR_CAPTURE_SRC}/hip_capture_scrub.cpp) - target_sources(hrr-unit-tests PRIVATE - hrr_scrub_test.cc ${HRR_CAPTURE_SRC}/hip_capture_scrub.cpp) + target_sources(hrr-unit-tests PRIVATE ${HRR_CAPTURE_SRC}/hip_capture_scrub.cpp) target_include_directories(hrr-unit-tests PRIVATE ${HRR_CAPTURE_SRC}) + target_compile_definitions(hrr-unit-tests PRIVATE HRR_HAVE_CAPTURE_SCRUB=1) endif() # ../include makes the local harness headers (hrr_test_common.hh, diff --git a/projects/hrr/tests/unit/hrr_scrub_test.cc b/projects/hrr/tests/unit/hrr_scrub_test.cc index fdef7fec66d1..d83fbd4c3859 100644 --- a/projects/hrr/tests/unit/hrr_scrub_test.cc +++ b/projects/hrr/tests/unit/hrr_scrub_test.cc @@ -15,6 +15,8 @@ */ #include "hrr_test_common.hh" + +#ifdef HRR_HAVE_CAPTURE_SCRUB #include "hip_capture_scrub.h" #include @@ -173,6 +175,15 @@ HRR_TEST_CASE(Unit_HRR_Scrub_RefusesOutOfBoundsTables) { } } +#else +// projects/hrr was checked out without the CLR sources (tests/unit/CMakeLists.txt). +#define HRR_SCRUB_NOT_BUILT(name) \ + HRR_TEST_CASE(name) { HRR_SKIP("built without CLR's hip_capture_scrub.cpp"); } +HRR_SCRUB_NOT_BUILT(Unit_HRR_Scrub_ZeroesAndRetypesBuildInfo) +HRR_SCRUB_NOT_BUILT(Unit_HRR_Scrub_SkipsUnterminatedName) +HRR_SCRUB_NOT_BUILT(Unit_HRR_Scrub_RefusesOutOfBoundsTables) +#endif + /** * @} */ diff --git a/projects/hrr/tools/tests/test_skill_api_ids.py b/projects/hrr/tools/tests/test_skill_api_ids.py new file mode 100644 index 000000000000..acb9ac94bf8d --- /dev/null +++ b/projects/hrr/tools/tests/test_skill_api_ids.py @@ -0,0 +1,51 @@ +# Copyright (c) Advanced Micro Devices, Inc., or its affiliates. +# +# SPDX-License-Identifier: MIT +"""The enable-recording skill knows the archive format it inspects. + +`inspect_archive.py verify` finds the fat binaries in an archive by the +`__hipRegisterFatBinary` id of the archive's format version, from a table it +keeps itself, and leaves them unscanned for a version it does not know. A +format bump that misses that table makes every new archive skip the scan +without any error, so the table has to map the current HRR_VERSION to the id +hrr_api_args.h gives that API. + +Run: python3 -m unittest discover -s projects/hrr/tools/tests +""" + +import importlib.util +import re +import sys +import unittest +from pathlib import Path + +HRR = Path(__file__).resolve().parents[2] +API_ARGS = HRR / "include" / "hrr" / "hrr_api_args.h" +INSPECT_ARCHIVE = HRR / "skills" / "enable-recording" / "scripts" / "inspect_archive.py" + + +def load_inspect_archive(): + spec = importlib.util.spec_from_file_location("inspect_archive", INSPECT_ARCHIVE) + module = importlib.util.module_from_spec(spec) + sys.modules[spec.name] = module # its dataclasses look the module up + spec.loader.exec_module(module) + return module + + +class SkillApiIdsTest(unittest.TestCase): + def test_fat_binary_id_covers_the_current_format_version(self): + text = API_ARGS.read_text(encoding="utf-8") + version = int(re.search(r"#define HRR_VERSION \(\(uint16_t\)(\d+)u\)", text).group(1)) + api_id = int(re.search(r"\bHRR_API_HIPREGISTERFATBINARY = (\d+),", text).group(1)) + table = load_inspect_archive().REGISTER_FAT_BINARY + self.assertEqual( + table.get(version), + api_id, + f"hrr_api_args.h is format v{version} with __hipRegisterFatBinary = {api_id}; " + "add that to REGISTER_FAT_BINARY in " + "skills/enable-recording/scripts/inspect_archive.py", + ) + + +if __name__ == "__main__": + unittest.main()