diff --git a/projects/clr/hipamd/src/hip_code_object.cpp b/projects/clr/hipamd/src/hip_code_object.cpp index a7974da74d0c..4b29d9f3e48a 100644 --- a/projects/clr/hipamd/src/hip_code_object.cpp +++ b/projects/clr/hipamd/src/hip_code_object.cpp @@ -725,10 +725,10 @@ hipError_t StatCO::InitManagedVarDevicePtr(int deviceId) { } // ================================================================================================ -void StatCO::ForEachFatBinaryBlob(void (*cb)(const void*)) const { +void StatCO::ForEachFatBinary(void (*cb)(FatBinaryInfo*)) { std::scoped_lock lock(sclock_); - for (const auto& [data, _] : modules_) { - cb(data); + for (const auto& [_, fat_binary_info] : modules_) { + cb(fat_binary_info); } } diff --git a/projects/clr/hipamd/src/hip_code_object.hpp b/projects/clr/hipamd/src/hip_code_object.hpp index eaf79ea32d02..a130b0c702aa 100644 --- a/projects/clr/hipamd/src/hip_code_object.hpp +++ b/projects/clr/hipamd/src/hip_code_object.hpp @@ -197,8 +197,9 @@ class StatCO : public CodeObject { // pointer is resolvable through MemObjMap before any HIP API can observe it. hipError_t EnsureAllManagedVarsAreMapped(); - // Iterate all registered fat binary data pointers — for HRR capture post-registration sweep. - void ForEachFatBinaryBlob(void (*cb)(const void*)) const; + // Iterate all registered fat binaries' FatBinaryInfo, null until HIP digests + // it; for HRR capture post-registration sweep. + void ForEachFatBinary(void (*cb)(FatBinaryInfo*)); // Iterate all registered __device__ globals as (host shadow address, symbol // name, size, device address) — the same post-registration sweep for HRR, diff --git a/projects/clr/hipamd/src/hrr/hip_capture.cpp b/projects/clr/hipamd/src/hrr/hip_capture.cpp index 074e79ba2c75..f826bff4c8cd 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture.cpp @@ -31,7 +31,6 @@ #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 @@ -49,6 +48,7 @@ #include "platform/kernel.hpp" // amd::KernelSignature #include "opencl/amdocl/cl_kernel.h" // T_POINTER enum #include "os/os.hpp" // amd::Os::installExceptionHandlers() +#include "elf/elf.hpp" // amd::Elf::getElfSize() // Fat binary format structs (ClangOffloadBundleUncompressedHeader, etc.) #include "../hip_code_object.hpp" @@ -57,6 +57,7 @@ #include #include #include +#include #include #include #include @@ -67,10 +68,18 @@ #include #include #include +#include #include #include #include +#if defined(__linux__) +#include // pipe2() +#include +#elif defined(_WIN32) +#include // VirtualQuery() +#endif + static std::once_flag g_hrr_atexit_once; @@ -1119,18 +1128,18 @@ static std::mutex g_prog_hash_mu; static std::unordered_map g_prog_hash; static std::atomic g_span_no_device_warned{false}; -// Fetch a program's device binary as (pointer, length) without dereferencing -// the bytes (the call itself still goes through binary_'s inserting -// operator[] — see above). +// Fetch a program's device binary for device dev as (pointer, length) without +// dereferencing the bytes (the call itself still goes through binary_'s +// inserting operator[] — see above). // Caller must hold g_prog_hash_mu. -// Returns false when there is no current device (or the index is out of -// range): that is not "the program has no binary", so callers must not seed -// g_prog_hash with a {0,0} / (nullptr,0) entry. -static bool program_binary_span_locked(amd::Program* prog, const uint8_t*& data, size_t& size) { +// Returns false when dev is not a device (no current device gives -1): that +// is not "the program has no binary", so callers must not seed g_prog_hash +// with a {0,0} / (nullptr,0) entry. +static bool program_binary_span_locked(amd::Program* prog, int dev, const uint8_t*& data, + size_t& size) { data = nullptr; size = 0; if (!prog) return false; - const int dev = hip::ihipGetDevice(); if (dev < 0 || static_cast(dev) >= hip::g_devices.size()) { if (!g_span_no_device_warned.exchange(true)) { LogPrintfWarning("[HRR capture] program %p: no current HIP device " @@ -1146,9 +1155,10 @@ static bool program_binary_span_locked(amd::Program* prog, const uint8_t*& data, return true; } -static bool program_binary_span(amd::Program* prog, const uint8_t*& data, size_t& size) { +static bool program_binary_span(amd::Program* prog, int dev, const uint8_t*& data, + size_t& size) { std::lock_guard lk(g_prog_hash_mu); - return program_binary_span_locked(prog, data, size); + return program_binary_span_locked(prog, dev, data, size); } static void remember_program_hash(const amd::Program* prog, const void* image, size_t size, @@ -1164,6 +1174,76 @@ static void forget_program_hash(const amd::Program* prog) { g_prog_hash.erase(prog); } +// How many of the n bytes at p can be read, found page by page without reading +// them, or nullopt when the probe cannot run. On Linux, write() of a page's +// first byte into a pipe fails with EFAULT where a plain read would fault, and +// no sandbox refuses write(); on Windows, VirtualQuery says whether the pages +// are committed and readable. +static std::optional readable_extent(const void* p, size_t n) { + const uintptr_t start = reinterpret_cast(p); + size_t ok = 0; +#if defined(__linux__) + const uintptr_t page = amd::Os::pageSize(); + int fds[2]; + if (pipe2(fds, O_CLOEXEC | O_NONBLOCK) != 0) return std::nullopt; + while (ok < n) { + char byte; + if (write(fds[1], reinterpret_cast(start + ok), 1) != 1) { + if (errno != EFAULT) ok = 0; + break; + } + (void)!read(fds[0], &byte, 1); // keep the pipe empty + ok = std::min(n, static_cast(((start + ok) | (page - 1)) + 1 - start)); + } + close(fds[0]); + close(fds[1]); +#elif defined(_WIN32) + constexpr DWORD kReadable = PAGE_READONLY | PAGE_READWRITE | PAGE_WRITECOPY | + PAGE_EXECUTE_READ | PAGE_EXECUTE_READWRITE | + PAGE_EXECUTE_WRITECOPY; + while (ok < n) { + MEMORY_BASIC_INFORMATION mbi; + if (!VirtualQuery(reinterpret_cast(start + ok), &mbi, sizeof(mbi)) || + mbi.State != MEM_COMMIT || !(mbi.Protect & kReadable) || (mbi.Protect & PAGE_GUARD)) + break; + const uintptr_t end = reinterpret_cast(mbi.BaseAddress) + mbi.RegionSize; + ok = std::min(n, static_cast(end - start)); + } +#endif + return ok; +} + +// The extent of the ELF image at p, of at most n bytes, as its sections and +// segments reach, 0 if it is not well formed or any of that cannot be read, or +// nullopt if readable_extent() cannot run. Probes the header, then the header +// tables, then the extent they give, so a declared size nothing bounds is never +// walked page by page. +static std::optional readable_elf_size(const uint8_t* p, size_t n) { + if (n < sizeof(Elf64_Ehdr)) return 0; + std::optional readable = readable_extent(p, sizeof(Elf64_Ehdr)); + if (!readable) return std::nullopt; + if (*readable < sizeof(Elf64_Ehdr)) return 0; + Elf64_Ehdr eh; + std::memcpy(&eh, p, sizeof(eh)); + uint64_t tables = sizeof(Elf64_Ehdr); + const uint64_t offs[2] = {eh.e_shoff, eh.e_phoff}; + const uint64_t lens[2] = {uint64_t(eh.e_shnum) * eh.e_shentsize, + uint64_t(eh.e_phnum) * eh.e_phentsize}; + for (int i = 0; i < 2; ++i) { + if (!lens[i]) continue; + if (offs[i] > n || lens[i] > n - offs[i]) return 0; + tables = std::max(tables, offs[i] + lens[i]); + } + readable = readable_extent(p, tables); + if (!readable) return std::nullopt; + if (*readable < tables) return 0; + const uint64_t size = amd::Elf::getElfSize(p, n); + if (!size) return 0; + readable = readable_extent(p, size); + if (!readable) return std::nullopt; + return *readable >= size ? size : 0; +} + // Hash an already-read device binary span, refusing anything that does not // look like a loadable image. Returns {0,0} when there is nothing safe to // record. Takes the span rather than re-reading it so callers that already @@ -1180,9 +1260,40 @@ static hrr_cap::Hash128 hash_program_image(const amd::Program* prog, const uint8 (const void*)prog, (unsigned long long)size, (const void*)data); return {0, 0}; } + // The span's length is not a readable extent: a code object HIP took from a + // registered fat binary carries the size its bundle entry declares, and + // nothing bounds that, nor the ELF's own headers, which can declare the same + // size. So hash no further than its sections and segments reach, and refuse + // one whose headers reach past the memory that can be read. + if (amd::Elf::isElfMagic(reinterpret_cast(data))) { + const std::optional elf_size = readable_elf_size(data, size); + if (!elf_size) { + LogPrintfWarning("[HRR capture] program %p: could not probe whether the ELF at %p" + " is readable; not recording a code object for it", + (const void*)prog, (const void*)data); + return {0, 0}; + } + if (*elf_size == 0) { + LogPrintfWarning("[HRR capture] program %p: ELF at %p is malformed or reaches past" + " readable memory within its %llu bytes; not recording a code" + " object for it", + (const void*)prog, (const void*)data, (unsigned long long)size); + return {0, 0}; + } + size = static_cast(*elf_size); + } return hrr_cap::writer::write_code_object(data, size); } +// hash_program_image(), and seed the launch-path hash cache (g_prog_hash) with +// the result for this span. +static hrr_cap::Hash128 record_program_image(const amd::Program* prog, const uint8_t* image, + size_t size) { + const hrr_cap::Hash128 h = hash_program_image(prog, image, size); + remember_program_hash(prog, image, size, h); + return h; +} + // Helper: get the actual device binary from a successfully loaded hipModule_t. // The module may come from an in-memory image or from a file (possibly a // bundle); either way the runtime stores the extracted device ELF in the @@ -1193,10 +1304,8 @@ static hrr_cap::Hash128 write_module_code_object(hipModule_t module) { if (!prog) return {0, 0}; const uint8_t* image = nullptr; size_t size = 0; - if (!program_binary_span(prog, image, size)) return {0, 0}; - hrr_cap::Hash128 h = hash_program_image(prog, image, size); - remember_program_hash(prog, image, size, h); - return h; + if (!program_binary_span(prog, hip::ihipGetDevice(), image, size)) return {0, 0}; + return record_program_image(prog, image, size); } static hrr_cap::Hash128 kernel_code_object_hash(amd::Kernel* kernel) { @@ -1211,7 +1320,7 @@ static hrr_cap::Hash128 kernel_code_object_hash(amd::Kernel* kernel) { // lookup in costs nothing since we hold the lock anyway. { std::lock_guard lk(g_prog_hash_mu); - got_span = program_binary_span_locked(prog, image, size); + got_span = program_binary_span_locked(prog, hip::ihipGetDevice(), image, size); if (got_span) { auto it = g_prog_hash.find(prog); if (it != g_prog_hash.end() && it->second.image == image && it->second.size == size) @@ -1221,9 +1330,7 @@ static hrr_cap::Hash128 kernel_code_object_hash(amd::Kernel* kernel) { if (!got_span) return {0, 0}; // Miss. Hash outside the lock — write_code_object takes the writer's own // locks, and hashing the same image twice is harmless (content-addressed). - hrr_cap::Hash128 h = hash_program_image(prog, image, size); - remember_program_hash(prog, image, size, h); - return h; + return record_program_image(prog, image, size); } hipError_t capture_hipModuleLoadData(hipModule_t* module, const void* image) { @@ -1624,121 +1731,57 @@ hipError_t capture_hipLaunchByPtr(const void* func) { } // --------------------------------------------------------------------------- -// Fat binary registration — record the binary blob +// Fat binary registration — record the code object HIP took from it // -// We capture the clang offload bundle blob so the replay can load it via -// hipModuleLoadData, making all kernel names resolvable. +// The bundle itself is not stored. HIP is handed a bare pointer and reads only +// the header and the code object it picks for each GPU, with a size COMGR +// gives it; a whole-bundle size could only come from the header, and nothing +// bounds that. Replay resolves kernels and globals in the code objects the +// archive holds, so the register event names the code object, not a blob. // --------------------------------------------------------------------------- -// Byte size of a clang offload bundle, or 0 when the header is not one or fails -// the checks HIP applies to it. HIP receives a registered fat binary as a bare -// pointer, so it knows no image bound (image_bound in hip_fatbin.cpp) and these -// checks are all that stands between the header and the read. -static size_t compute_bundle_size(const void* blob) { - if (!blob) return 0; - const char* p = static_cast(blob); - - // Compressed format: magic "CCOB", header contains totalSize at byte 8. - if (std::memcmp(p, hip::symbols::kOffloadBundleCompressedMagicStr, - hip::symbols::kOffloadBundleCompressedMagicStrSize - 1) == 0) { - constexpr size_t kHeaderSize = - offsetof(hip::symbols::ClangOffloadBundleCompressedHeader, compressedBinarydesc); - const auto* hdr = static_cast(blob); - if (hdr->totalSize < kHeaderSize) { - LogPrintfWarning( - "[HRR capture] Fat binary at %p not recorded: compressed bundle " - "totalSize %u is smaller than its %zu-byte header", - blob, static_cast(hdr->totalSize), kHeaderSize); - return 0; - } - return static_cast(hdr->totalSize); - } - - // Uncompressed format: magic "__CLANG_OFFLOAD_BUNDLE__" + numOfCodeObjects + entries. - if (std::memcmp(p, hip::symbols::kOffloadBundleUncompressedMagicStr, - hip::symbols::kOffloadBundleUncompressedMagicStrSize - 1) != 0) { - return 0; // Unknown format - } - const auto* hdr = static_cast(blob); - uint64_t n = hdr->numOfCodeObjects; - if (n == 0) return 0; - - // HIP hands COMGR only the first 4096 bytes, and COMGR reads the entry table from - // that slice and fails if a read runs off its end. __hipRegisterFatBinary passes - // no length, so this bound is also what keeps the walk below in range. - constexpr size_t kEntryTableLimit = 4096; - constexpr size_t kEntryHeaderSize = offsetof(hip::symbols::ClangOffloadBundleInfo, bundleEntryId); - size_t pos = offsetof(hip::symbols::ClangOffloadBundleUncompressedHeader, desc); - uint64_t end = 0; - for (uint64_t i = 0; i < n; i++) { - const auto* entry = reinterpret_cast(p + pos); - if (kEntryTableLimit - pos < kEntryHeaderSize || - entry->bundleEntryIdSize > kEntryTableLimit - pos - kEntryHeaderSize) { - LogPrintfWarning( - "[HRR capture] Fat binary at %p not recorded: bundle entry %llu ends " - "past the first %zu bytes, where HIP looks for it", - blob, static_cast(i), kEntryTableLimit); - return 0; - } - pos += kEntryHeaderSize + static_cast(entry->bundleEntryIdSize); - // HIP makes this check only for the code objects it picks for this host's GPUs. - // Every entry counts towards the end here, so capture checks them all. - const uint64_t entry_end = entry->offset + entry->size; - if (entry_end < entry->offset) { - LogPrintfWarning( - "[HRR capture] Fat binary at %p not recorded: bundle entry %llu has " - "an offset and size that overflow", - blob, static_cast(i)); - return 0; - } - if (entry_end > end) end = entry_end; - } - static_assert(sizeof(size_t) >= sizeof(uint64_t), "a 64-bit bundle size must fit in size_t"); - return static_cast(end); +// Record the code object HIP extracted from a digested fat binary for the +// current device, or device 0 before one is current, the device the global +// variable sweep resolves on. A fat binary HIP has not digested yet has +// nothing loaded; a launch from it records its code object then. +// +// Returns the code object's hash, which the register event names. The size +// HIP holds is the bundle entry's declared one, so only an ELF, whose headers +// bound it (see hash_program_image), is recorded. +static hrr_cap::Hash128 record_fat_binary_code_object(hip::FatBinaryInfo* fbi) { + if (!fbi) return {0, 0}; + const int dev = std::max(hip::ihipGetDevice(), 0); + // GetProgram() aborts on a device it does not have. + if (static_cast(dev) >= hip::g_devices.size()) return {0, 0}; + amd::Program* prog = fbi->GetProgram(dev); + const uint8_t* image = nullptr; + size_t size = 0; + if (!program_binary_span(prog, dev, image, size) || !image || size < sizeof(Elf64_Ehdr) || + !amd::Elf::isElfMagic(reinterpret_cast(image))) + return {0, 0}; + return record_program_image(prog, image, size); } -// 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), - 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(), /*capped=*/false); +// The register event names the code object in its blob hash, with a blob +// size of 0, which tells it from a fat binary blob. +static void write_register_fat_binary_event(uint64_t ret, hrr_cap::Hash128 co) { + hrr_args___hipRegisterFatBinary a{}; + a.ret = ret; + a.blob_hash_lo = co.lo; + a.blob_hash_hi = co.hi; + hrr_cap::writer::write_event_raw(HRR_API_HIPREGISTERFATBINARY, &a.hdr, sizeof(a)); } 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. - // data is a __CudaFatBinaryWrapper* { magic, version, binary, dummy }. - // Capture the fat binary blob (binary field) so replay can load it via hipModuleLoadData. - struct __HRRFatBinaryWrapper { uint32_t magic; uint32_t version; const void* binary; const void* dummy; }; - const auto* wrapper = static_cast(data); - const void* blob = (wrapper && (wrapper->magic == 0x48495046u /*HIPF*/ || - wrapper->magic == 0x4B504948u /*HIPK*/)) - ? wrapper->binary : nullptr; // PlatformState is initialized before this shim is installed, so HIP has already - // digested the bundle, and no handle means it could not load it on this host: a bad - // header or, more often, no code object for any GPU here. Then no launch from it can - // succeed and be recorded, so replay never needs the blob. - if (blob && !r) { - LogPrintfWarning( - "[HRR capture] Fat binary at %p not recorded: HIP did not register it " - "(no code object for any GPU here, or a bad header)", - blob); - blob = nullptr; - } - size_t blob_size = blob ? compute_bundle_size(blob) : 0; - - hrr_args___hipRegisterFatBinary a{}; - a.ret = reinterpret_cast(r); - a.blob_size = static_cast(blob_size); - if (blob && blob_size > 0) { - auto h = write_fat_binary_blob(blob, 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)); + // digested the bundle and r, when set, is its FatBinaryInfo**. + const hrr_cap::Hash128 co = + r ? record_fat_binary_code_object(*reinterpret_cast(r)) + : hrr_cap::Hash128{0, 0}; + write_register_fat_binary_event(reinterpret_cast(r), co); return r; } @@ -2863,23 +2906,12 @@ void hip_capture_install_compiler_table(const HipCompilerDispatchTable& shims) { // Init / Shutdown // --------------------------------------------------------------------------- -// Record a single fat binary blob as a HRR_API_HIPREGISTERFATBINARY event. -// blob_ptr is the fbwrapper->binary pointer (the actual clang offload bundle). -// These were registered before HIP initialized and are not digested yet, so -// unlike the live shim there is no HIP verdict to check, only the header, and a -// bundle with no code object for this host is still recorded. -static void record_fat_binary_blob(const void* blob_ptr) { - if (!blob_ptr) return; - size_t blob_size = compute_bundle_size(blob_ptr); - if (blob_size == 0) return; - - hrr_args___hipRegisterFatBinary a{}; - a.ret = 0; // handle not meaningful at init time - a.blob_size = static_cast(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)); +// Record a fat binary registered before capture started as a +// HRR_API_HIPREGISTERFATBINARY event, with the code object HIP has loaded from +// it so far (fbi is null until HIP digests it). +static void record_fat_binary(hip::FatBinaryInfo* fbi) { + // ret 0: the handle is not meaningful at init time. + write_register_fat_binary_event(0, record_fat_binary_code_object(fbi)); } // Record one registered __device__ global as a HRR_API_HIPREGISTERVAR event. @@ -3010,14 +3042,13 @@ void hip_capture_init() { // the compiler dispatch table is fully populated. hip_capture_build_compiler_table(); - // Retroactively record fat binaries that fired before our shims were live. - // __hipRegisterFatBinary fires at app static-init, before hip_capture_init(). - hip::PlatformState::Instance().StatCO().ForEachFatBinaryBlob(record_fat_binary_blob); - - // And the __device__ globals registered against them, for the same reason - // and in this order: the fat binaries have to be in the archive before the - // symbols that live inside them. + // Retroactively record the __device__ globals and fat binaries registered + // before our shims were live: both fire at app static-init, before + // hip_capture_init(). Globals first, because resolving one makes HIP load + // the code object holding it, which the fat binary sweep then records. + // Replay loads every code object before it replays any event. hip::PlatformState::Instance().StatCO().ForEachGlobalVar(record_registered_var); + hip::PlatformState::Instance().StatCO().ForEachFatBinary(record_fat_binary); std::call_once(g_hrr_atexit_once, [] { std::atexit(hip_capture_shutdown); }); @@ -3033,7 +3064,7 @@ void hip_capture_init() { #else // HRR capture is disabled // below to avoid -Wunused-function / -Wunused-variable - (void)&record_fat_binary_blob; + (void)&record_fat_binary; (void)&record_registered_var; (void)&hrr_install_clr_exception_handler; (void)&g_hrr_atexit_once; diff --git a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp index a67f41561c5c..fc385c77bad6 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp +++ b/projects/clr/hipamd/src/hrr/hip_capture_writer.cpp @@ -1024,9 +1024,8 @@ static uint64_t file_bytes(uint64_t len) { // still opens an archive of its own, held to the same limit. // HIP_HRR_CAPTURE_MAX_BLOB_MB caps one blob: a larger copy is not stored, its // event records no blob, the archive is marked incomplete and capture goes on. -// Code objects and the fat binaries HIP registers are not held to it; a -// hipLinkAddData image is. 0, the default, sets no limit. Both apply with or -// without the reserve. +// Code objects are not held to it; a hipLinkAddData image is. 0, the default, +// sets no limit. Both apply with or without the reserve. // // A write over the archive limit is refused and capture stopped under the same // hold of g_file_mu, which flush() takes to decide on the trailer: an archive @@ -1811,13 +1810,13 @@ class BlobClaim { // write_blob // --------------------------------------------------------------------------- -Hash128 write_blob(const void* data, size_t len, bool capped) { +Hash128 write_blob(const void* data, size_t len) { reopen_after_fork(); { std::lock_guard lk(g_file_mu); if (g_events_fd < 0 || g_events_finalized) return {}; // not open, or past the trailer } - if (capped && g_max_blob_bytes != 0 && len > g_max_blob_bytes) { + if (g_max_blob_bytes != 0 && len > g_max_blob_bytes) { skip_oversized_blob(len); return {}; } diff --git a/projects/clr/hipamd/src/hrr/hip_capture_writer.h b/projects/clr/hipamd/src/hrr/hip_capture_writer.h index a27483068d41..affc69f0ec2d 100644 --- a/projects/clr/hipamd/src/hrr/hip_capture_writer.h +++ b/projects/clr/hipamd/src/hrr/hip_capture_writer.h @@ -91,12 +91,10 @@ void note_unreplayable(const char* api, const char* reason); // Write a buffer as a content-addressed blob. Returns hash, or {} when the // writer is not open, flush() has written the trailer, or capture has stopped -// for lack of space or at the size limit. A `capped` buffer over -// HIP_HRR_CAPTURE_MAX_BLOB_MB is not written either: {} is returned and the -// capture marked incomplete. Fat binaries are not capped: replay loads them to -// run their kernels and resolve their device globals. +// for lack of space or at the size limit. A buffer over HIP_HRR_CAPTURE_MAX_BLOB_MB +// is not written either: {} is returned and the capture marked incomplete. // Thread-safe. Skips the write when the same bytes are already in this archive. -Hash128 write_blob(const void* data, size_t len, bool capped = true); +Hash128 write_blob(const void* data, size_t len); // Write a code object (.hsaco) blob. Returns hash, or {} as write_blob() does. Hash128 write_code_object(const void* image, size_t image_size); diff --git a/projects/hrr/DESIGN.md b/projects/hrr/DESIGN.md index fe4bad9506e1..d85e6da19a66 100644 --- a/projects/hrr/DESIGN.md +++ b/projects/hrr/DESIGN.md @@ -52,7 +52,7 @@ through one of these shims, which records a fixed-size `hrr_event_header` (32 by type, sequence ID, timestamp, thread ID) followed by a variable-length payload of serialised arguments into a streaming `events.bin` file. -A subset of APIs requires hand-written shims (`MANUAL_CAPTURE_APIS`): kernel launches (4 variants), host-to-device memcpy (4 variants), module load (3 variants), `__hipRegisterFatBinary`, host memory registration (`hipHostRegister`, `hipHostUnregister`), all `hipMemcpy3D` variants (4), array creation, VMM, stream/memory attribute APIs, and others where struct pointer serialisation requires custom handling. All others are code-generated by `gen_hrr_api_args.py` from `hip_api_trace.hpp`. Fat binaries registered before shim install (C++ static-init) are recovered retroactively via `ForEachFatBinaryBlob` at init. +A subset of APIs requires hand-written shims (`MANUAL_CAPTURE_APIS`): kernel launches (4 variants), host-to-device memcpy (4 variants), module load (3 variants), `__hipRegisterFatBinary`, host memory registration (`hipHostRegister`, `hipHostUnregister`), all `hipMemcpy3D` variants (4), array creation, VMM, stream/memory attribute APIs, and others where struct pointer serialisation requires custom handling. All others are code-generated by `gen_hrr_api_args.py` from `hip_api_trace.hpp`. Fat binaries registered before shim install (C++ static-init) are recovered retroactively via `ForEachFatBinary` at init. H2D memcpy captures the host source buffer as a content-addressed blob (FNV-1a-128 hash). D2H memcpy captures the expected output blob after a forced stream sync. Kernel argument pointers are detected via `kernel->signature()` (in-tree access to `amd::KernelSignature`) — no ELF parsing needed. @@ -189,7 +189,7 @@ capture.hrr/ manifest.json { pid, parent_pid, complete, event_count, blob_count } active names this process instance: present while its capture is on (see Transport) blobs/<2hex>/ content-addressed host buffers keyed by FNV-1a-128 hash - code_objects/ .hsaco ELFs (unused in current fat-binary path) + code_objects/ .hsaco ELFs: loaded modules and the code object HIP takes from each fat binary regions/*.hrrr external region annotations (optional; written by producers outside the runtime, never by capture — see below) pid-.prev-/ an earlier process's archive with the same pid, moved @@ -508,7 +508,7 @@ manually: `const uint8_t* p = pl + sizeof(hrr_event_header);` | API | Extra fields | |-----|-------------| -| `__hipRegisterFatBinary` | `blob_hash_lo u64`, `blob_hash_hi u64`, `blob_size u64` | +| `__hipRegisterFatBinary` | `blob_hash_lo u64`, `blob_hash_hi u64`, `blob_size u64` (0 when the hash names the code object HIP loaded) | | `hipModuleLoadData`, `hipModuleLoadDataEx`, `hipModuleLoad` | `co_hash_lo u64`, `co_hash_hi u64`, `module_id u32` | | `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync` | `blob_hash_lo u64`, `blob_hash_hi u64` | | `hipHostRegister` | `blob_hash_lo u64`, `blob_hash_hi u64` | @@ -794,9 +794,15 @@ not its type, so the type is not compared. `__hipRegisterFatBinary` fires at C++ static-init time before shims are installed. At init, all already-registered fat binaries are recovered retroactively via -`ForEachFatBinaryBlob`. At playback, each captured blob is loaded with -`hipModuleLoadData`; the resulting `hipModule_t` is stored in `module_map` keyed by -hash and scanned for kernel name lookups. +`ForEachFatBinary`. Capture does not store the bundle: its size could only come +from the header, which nothing bounds. It records the code object HIP extracted +from the bundle for the current device, bounded by the ELF's own headers and by the +memory that can be read, since the size HIP knows is the one the bundle entry declares +and the headers can declare as much; one whose headers reach further is not recorded. +The register event names that code object (blob hash set, blob size 0). Replay pre-loads every code object, +so kernels and `__device__` globals resolve as they do for module loads. A register +event with a blob size carries the fat binary itself as a blob, which playback loads +with `hipModuleLoadData`. ## Multi-Threaded Capture @@ -1212,7 +1218,7 @@ kernel. `__hipRegisterFatBinary` fires at C++ static-initialisation time, before `hip_capture_init()` installs shims. All fat binaries registered before the shim is active are missed by the normal capture path and recovered retroactively via -`PlatformState::Instance().StatCO().ForEachFatBinaryBlob(...)` at init time. +`PlatformState::Instance().StatCO().ForEachFatBinary(...)` at init time. If an application loads additional shared libraries _after_ HIP init that register their own fat binaries, those will be captured normally through the shim. The edge diff --git a/projects/hrr/PRIVACY.md b/projects/hrr/PRIVACY.md index 8b6538f0e4c0..4ec51133f217 100644 --- a/projects/hrr/PRIVACY.md +++ b/projects/hrr/PRIVACY.md @@ -57,11 +57,10 @@ not make. | Host buffers copied to or from the device, byte for byte, unencrypted. For a serving stack this includes prompts and generated output as token ids, which the model's public tokeniser reverses, and the model weights | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | | Host memory registered with `hipHostRegister`: the whole range, byte for byte, as it was at registration, whether or not it is ever copied | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | | Kernel argument bytes | `events.bin` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | -| Code objects the process loaded (the application's own GPU binary code) | `code_objects/`, `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | +| Code objects the process loaded (the application's own GPU binary code); of a fat binary, the code object HIP loads, which its `__hipRegisterFatBinary` record names with a blob size of 0. A record with a blob size names a whole bundle in `blobs/`, with targets HIP never loads | `code_objects/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880); [AIRUNTIME-2858](https://amd-hub.atlassian.net/browse/AIRUNTIME-2858): of a fat binary, only the code object HIP loads | | Linker input buffers passed to `hipLinkAddData`, every byte, with their names, including inputs the call rejected and inputs never linked or loaded. A SPIR-V or LLVM bitcode input keeps its compiler identification and, built with `-g`, its source and build directories; capture does not scrub them | `blobs/` | contents | [AIRUNTIME-2880](https://amd-hub.atlassian.net/browse/AIRUNTIME-2880) | -| Whole fat binary bundles, including targets HIP never loads | `blobs/` | contents | [AIRUNTIME-2858](https://amd-hub.atlassian.net/browse/AIRUNTIME-2858): only the code objects HIP loads | -| Debug and comment sections of code objects: absolute build paths, source file names, the exact compiler build | `code_objects/`, `blobs/` | contents | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): zeroed in code objects and uncompressed fat binaries; kept in compressed (`CCOB`) fat binaries | -| Strings in the code the GPU loads, such as the absolute source path a device `assert()` prints | `code_objects/`, `blobs/` | contents | — kept, the kernel uses them; `verify` counts code objects holding an absolute path | +| Debug and comment sections of code objects: absolute build paths, source file names, the exact compiler build | `code_objects/` | contents | [AIRUNTIME-2881](https://amd-hub.atlassian.net/browse/AIRUNTIME-2881): zeroed in every code object stored, including the one HIP takes from a compressed (`CCOB`) fat binary. A compressed bundle that a record with a blob size names in `blobs/` keeps them | +| Strings in the code the GPU loads, such as the absolute source path a device `assert()` prints | `code_objects/` | contents | — kept, the kernel uses them; `verify` counts code objects holding an absolute path | | Framework allocator layout, when a region producer ran beside capture | `regions/` | both | — kept, written only by an opt-in producer | | Marker that capture is running: on Linux the boot id and the process start time in clock ticks since boot; on Windows the process creation date and time. A normal exit removes it; a process that crashes or is killed leaves it | `active` | both | — kept, region producers read it | diff --git a/projects/hrr/README.md b/projects/hrr/README.md index 63b50b423eb0..4956dc112cc9 100644 --- a/projects/hrr/README.md +++ b/projects/hrr/README.md @@ -94,7 +94,7 @@ User-facing capture, replay, and validation knobs. Implementation details can be | `HIP_HRR_DEBUG_ARGS` | off | Dump every captured kernel arg to the log (debug / provenance). Only a Debug build of the HIP runtime (no `NDEBUG`) reads it; a release build ignores it. Either way capture says on stderr whether it reads it, whatever the log level | | `HIP_HRR_HOST_SNAPSHOTS` | on | Record the pinned host memory each kernel reads before the launch, so replay can write it back. `HIP_HRR_HOST_SNAPSHOT_MAX_MB` (default `64`) caps the allocation size. The archive then holds that host memory (see DESIGN.md, Pinned Host Snapshots) | | `HIP_HRR_CAPTURE_MAX_MB` | 0 (no limit) | Largest size, in MiB, of each process's archive (`pid-/`), manifests aside. The write that would pass it is not made: capture stops with a `[HRR capture] Capture stopped` message on stderr, the archive is marked incomplete, and the application runs on. A process it forks afterwards still gets an archive of its own | -| `HIP_HRR_CAPTURE_MAX_BLOB_MB` | 0 (no limit) | Largest copy, in MiB, stored as a blob. A larger one is recorded without its contents, the archive is marked incomplete with one message on stderr, and capture goes on. Replay refuses by name a host-to-device copy recorded so, whether `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync`, `hipMemcpyWithStream`, a 2D or 3D copy or a to-symbol graph copy, a `hipHostRegister`, and a kernel launch outside a stream capture whose pinned host snapshot was left out. Code objects and the fat binaries HIP registers are always stored; a `hipLinkAddData` image is not, and replay refuses its link | +| `HIP_HRR_CAPTURE_MAX_BLOB_MB` | 0 (no limit) | Largest copy, in MiB, stored as a blob. A larger one is recorded without its contents, the archive is marked incomplete with one message on stderr, and capture goes on. Replay refuses by name a host-to-device copy recorded so, whether `hipMemcpy`, `hipMemcpyAsync`, `hipMemcpyHtoD`, `hipMemcpyHtoDAsync`, `hipMemcpyWithStream`, a 2D or 3D copy or a to-symbol graph copy, a `hipHostRegister`, and a kernel launch outside a stream capture whose pinned host snapshot was left out. Code objects are always stored; a `hipLinkAddData` image is not, and replay refuses its link | On Linux the archive is readable only by the user who captured it (directories 0700, files 0600). What an archive contains, which of it each capture mode records, and why replaying an archive means running its GPU code: [PRIVACY.md](PRIVACY.md). If the archive cannot be set up, capture is disabled with a `[HRR capture] Capture disabled` message on stderr and the application runs normally. @@ -248,11 +248,12 @@ capture.hrr/ still resolve to the code object they came from. Their `.comment` and `.debug*` sections, which hold the compiler build and the build paths of debug information, are zeroed and retyped `SHT_NOBITS` before they are - stored, as are those of uncompressed fat binaries, so a debugger on a - replay sees no debug information. What the GPU loads is not touched: - kernel names stay, and so does a string a kernel keeps, such as the - source path a device `assert()` prints. Compressed (`CCOB`) fat binaries - and `hipLinkAddData` inputs are stored as they were passed. + stored, so a debugger on a replay sees no debug information. Of a fat + binary, compressed or not, only the code object HIP loads is stored, here + and scrubbed the same way. What the GPU loads is not touched: kernel + names stay, and so does a string a kernel keeps, such as the source path + a device `assert()` prints. `hipLinkAddData` inputs are stored as they + were passed. - **manifest.json** — counts, versions and device properties; the GPU's PCI address, UUID and LUID are written blank - **Complete: NO** — original run crashed before clean shutdown; reader still recovers complete events diff --git a/projects/hrr/include/hrr/hrr_api_args.h b/projects/hrr/include/hrr/hrr_api_args.h index a5de4f04a9d0..ca2f4f6542e8 100644 --- a/projects/hrr/include/hrr/hrr_api_args.h +++ b/projects/hrr/include/hrr/hrr_api_args.h @@ -210,7 +210,7 @@ typedef struct { typedef struct { hrr_event_header hdr; uint64_t ret; - uint64_t blob_hash_lo; /* FNV-1a-128 lo of fat binary */ + uint64_t blob_hash_lo; /* FNV-1a-128 lo of fat binary, or of its code object if blob_size is 0 */ uint64_t blob_hash_hi; /* FNV-1a-128 hi */ uint64_t blob_size; /* fat binary byte count */ } hrr_args___hipRegisterFatBinary; diff --git a/projects/hrr/playback/hip_playback.cpp b/projects/hrr/playback/hip_playback.cpp index ea8056b945ba..0c14cb630403 100644 --- a/projects/hrr/playback/hip_playback.cpp +++ b/projects/hrr/playback/hip_playback.cpp @@ -1759,6 +1759,7 @@ hipError_t playback___hipRegisterFatBinary(PlaybackContext& ctx, uint64_t blob_hash_hi = a->blob_hash_hi; if (!blob_hash_lo && !blob_hash_hi) return hipSuccess; // no blob — skip + // No blob: the hash names a code object, which run_pass() pre-loads. if (!a->blob_size) return hipSuccess; std::string hex = hrr::hash_hex(blob_hash_lo, blob_hash_hi); diff --git a/projects/hrr/skills/enable-recording/SKILL.md b/projects/hrr/skills/enable-recording/SKILL.md index fac5010c968a..7c41cfd933c3 100644 --- a/projects/hrr/skills/enable-recording/SKILL.md +++ b/projects/hrr/skills/enable-recording/SKILL.md @@ -130,13 +130,13 @@ 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. 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. +- **`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. Of a fat binary, only the code object HIP loads is stored, not the bundle with its other targets. Every code object stored has its `.comment` and `.debug*` sections zeroed and marked empty, which removes the compiler build and the build paths and source file names that debug information carries; this holds for code objects HIP took from compressed (`CCOB`) fat binaries too. 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`; `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 v12 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 v12 and it cannot tell which blobs are fat binaries. Both fat binary fields concern the bundles that registration records with a blob size name in `blobs/`; a record with a blob size of 0 names the code object HIP loaded, which is scanned under `code_objects/`); 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/tests/config/integration.yaml b/projects/hrr/tests/config/integration.yaml index 9a2c728a2198..5446d1b9e40c 100644 --- a/projects/hrr/tests/config/integration.yaml +++ b/projects/hrr/tests/config/integration.yaml @@ -159,6 +159,8 @@ cases: Unit_HRR_FailedMemcpy3D_Direct: {} Unit_HRR_FailedShimCallsNotRecorded: {} Unit_HRR_FailedShimCalls_Direct: {} + Unit_HRR_FatBinGlobalCopy: {} + Unit_HRR_FatBinGlobalCopy_Direct: {} Unit_HRR_FatBinHeaderChecks: {} Unit_HRR_FatBinHeader_Direct: {} Unit_HRR_FirstMallocRoundtrip: {} diff --git a/projects/hrr/tests/include/hrr_test_common.hh b/projects/hrr/tests/include/hrr_test_common.hh index eeb04b09c03e..798d1ad5efd9 100644 --- a/projects/hrr/tests/include/hrr_test_common.hh +++ b/projects/hrr/tests/include/hrr_test_common.hh @@ -329,9 +329,10 @@ inline void hrr_run_playback(const fs::path& cap_path, // must produce at least a few events (malloc, memcpy, kernel, free) — a value // of 5 is a conservative floor that would catch a totally empty capture. // Use a higher value for workloads known to emit many events (StressApis, etc.). -// require_d2h: if true (default), asserts that playback validated at least one -// D2H blob. Pass false for workloads that conditionally skip D2H (e.g. the -// texture workload on devices without image support). +// require_d2h: if true (default), asserts that the archive holds a blob and +// that playback validated at least one D2H blob. Pass false for workloads +// that conditionally skip D2H (e.g. the texture workload on devices without +// image support), whose archive may hold no blob at all. // --------------------------------------------------------------------------- inline void hrr_run_roundtrip(const std::string& direct_case, const fs::path& cap_path, @@ -351,7 +352,7 @@ inline void hrr_run_roundtrip(const std::string& direct_case, int bc = 0; for ([[maybe_unused]] const auto& _ : fs::recursive_directory_iterator(archive_path / "blobs")) ++bc; - INFO("Blob count: " << bc); REQUIRE(bc >= 1); + INFO("Blob count: " << bc); if (require_d2h) REQUIRE(bc >= 1); // Load the archive and assert a minimum event count. This catches generator // bugs that silently produce empty or near-empty archives while still writing diff --git a/projects/hrr/tests/integration/CMakeLists.txt b/projects/hrr/tests/integration/CMakeLists.txt index 3442a9090458..7cc2b4511db8 100644 --- a/projects/hrr/tests/integration/CMakeLists.txt +++ b/projects/hrr/tests/integration/CMakeLists.txt @@ -72,9 +72,10 @@ # It reaches the live dispatch table through rocprofiler-register, with # hrr_dispatch_tool.cc as the tool library. Linux only. # -# hrr_fatbin_header_test.cc covers which registered fat binaries capture stores: -# its hidden [.][hrr-direct] workload registers valid and altered bundle -# headers itself, before and after HIP initialises. It compiles its kernel with +# hrr_fatbin_header_test.cc covers what capture stores of a registered fat +# binary, the code object HIP loads and none of the bundle: its hidden +# [.][hrr-direct] workload registers valid and altered bundle headers itself, +# before and after HIP initialises. It compiles its kernel with # hipRTC, so it is a plain C++ translation unit. # # hrr_disk_space_test.cc covers how capture stops before the file system diff --git a/projects/hrr/tests/integration/hrr_fatbin_header_test.cc b/projects/hrr/tests/integration/hrr_fatbin_header_test.cc index aa3e8cde38d3..22193b79636d 100644 --- a/projects/hrr/tests/integration/hrr_fatbin_header_test.cc +++ b/projects/hrr/tests/integration/hrr_fatbin_header_test.cc @@ -8,29 +8,26 @@ * @addtogroup HRR HRR fat binary header checks * @{ * @ingroup HRRTest - * Tests for which registered fat binaries capture stores as blobs. + * Tests for what capture stores of a registered fat binary. * - * Capture stores each registered fat binary so replay can load it with - * hipModuleLoadData, and the byte count it stores comes from the bundle - * header. HIP checks that header before it loads the bundle, and capture - * applies the same checks: a bundle whose header fails them is not stored, - * and neither is a bundle HIP refused to register after it initialised. + * Capture stores none of the bundle: HIP is handed a bare pointer, so the + * bundle's extent is known only from its header, and HIP itself reads just the + * header and the code object it picks for this GPU. Capture records that code + * object, cut to the extent its ELF headers give, and the register event + * without a blob. A bundle whose header declares an entry far past the real + * image is captured without reading there. * * Two paths record fat binaries and both are exercised: * - registered before HIP initialises: swept from HIP's list in - * hip_capture_init, where only the header checks apply; - * - registered after: seen by the __hipRegisterFatBinary shim, which also - * has HIP's verdict. + * hip_capture_init, after HIP has loaded what it needed for globals; + * - registered after: seen by the __hipRegisterFatBinary shim, once HIP has + * digested the bundle. * The workload calls __hipRegisterFatBinary itself with bundles laid out in * memory the way clang-offload-bundler lays them out, valid and altered. * - * Every code object capture stores, in a fat binary or under code_objects/, - * has its .comment and .debug* sections zeroed: they hold the compiler build - * and the build paths, and the GPU never loads them. - * - * A fat binary is stored whatever its size, HIP_HRR_CAPTURE_MAX_BLOB_MB - * notwithstanding: replay loads it to run its kernels and resolve its device - * globals. + * Every code object capture stores under code_objects/ has its .comment and + * .debug* sections zeroed: they hold the compiler build and the build paths, + * and the GPU never loads them. */ #include "hrr_test_common.hh" @@ -39,15 +36,30 @@ #include #include +#include #include #include +#include #include +#include #include #include +#if !defined(_WIN32) +#include +#endif + // Exported by libamdhip64 and called by the module constructor clang emits for // every HIP translation unit; no header declares it. extern "C" void** __hipRegisterFatBinary(const void* data); +// These two take names as char* but only read them, so the calls below +// const_cast string literals to match the C signature. +extern "C" void __hipRegisterFunction(void** modules, const void* hostFunction, + char* deviceFunction, const char* deviceName, + unsigned int threadLimit, uint3* tid, uint3* bid, + dim3* blockDim, dim3* gridDim, int* wSize); +extern "C" void __hipRegisterVar(void** modules, void* var, char* hostVar, char* deviceVar, + int ext, size_t size, int constant, int global); #define HRR_HIPRTC_CHECK(expr) \ do { \ @@ -75,9 +87,8 @@ constexpr char kHostId[] = "host-x86_64-unknown-linux-gnu-"; constexpr char kTargetPrefix[] = "hipv4-amdgcn-amd-amdhsa--"; // A real target no runner has, so HIP finds no code object for its GPUs. constexpr char kForeignId[] = "hipv4-amdgcn-amd-amdhsa--gfx600"; -// The compressed header as HIP reads it (version 2): 24 bytes up to the data. -constexpr size_t kCompressedHeaderSize = 24; -constexpr uint32_t kHipfMagic = 0x48495046u; // "HIPF" +constexpr uint32_t kHipfMagic = 0x48495046u; // "HIPF" +constexpr uint64_t kTerabyte = uint64_t(1) << 40; struct BundleEntry { uint64_t offset; @@ -121,48 +132,35 @@ std::vector foreign_bundle(const char* marker) { return b; } -// A compressed header whose totalSize does not even cover the header. The -// hash field carries the tag, so the bytes a capture stored can be told apart. -std::vector short_compressed_bundle(uint64_t tag) { - std::vector b(kCompressedHeaderSize, 0); - std::memcpy(b.data(), "CCOB", 4); - const uint16_t version = 2, method = 1; - const uint32_t total_size = kCompressedHeaderSize - 1, uncompressed_size = 0; - std::memcpy(b.data() + 4, &version, 2); - std::memcpy(b.data() + 6, &method, 2); - std::memcpy(b.data() + 8, &total_size, 4); - std::memcpy(b.data() + 12, &uncompressed_size, 4); - put_u64(b, 16, tag); - return b; -} - -// One entry whose identifier runs past the first 4096 bytes. The entry it -// describes is in range, so only the table bound refuses it. -std::vector long_id_bundle(const char* marker) { - std::string id(kTableLimit, '\0'); - std::memcpy(&id[0], kForeignId, sizeof(kForeignId) - 1); - const size_t payload = kTableStart + kEntryHeaderSize + id.size(); - auto b = uncompressed_bundle({{payload, kMarkerPayload, id}}, payload + kMarkerPayload); - put_marker(b, payload, marker); - return b; -} - -// A host entry whose offset plus size wraps around, beside a code object for -// target_id. HIP never loads the host entry and does not check it. -std::vector host_overflow_bundle(const std::string& target_id, - const std::vector& payload) { - auto b = uncompressed_bundle({{1, UINT64_MAX, kHostId}, - {kPayloadOffset, payload.size(), target_id}}, - kPayloadOffset + payload.size()); - std::memcpy(b.data() + kPayloadOffset, payload.data(), payload.size()); - return b; +// The code object with its first section the loader never reads (PROGBITS, +// not allocated: .comment as lld emits it) declaring a terabyte from the start +// of the file. HIP looks only at the ELF header of a shared object it is given. +std::vector overstated_section(std::vector co) { + uint64_t shoff = 0, flags = 0; + uint16_t shnum = 0; + uint32_t type = 0; + std::memcpy(&shoff, co.data() + 0x28, sizeof(shoff)); + std::memcpy(&shnum, co.data() + 0x3c, sizeof(shnum)); + for (size_t sh = shoff; sh < shoff + size_t(shnum) * 64; sh += 64) { + std::memcpy(&type, co.data() + sh + 4, sizeof(type)); + std::memcpy(&flags, co.data() + sh + 8, sizeof(flags)); + if (type != 1 /* SHT_PROGBITS */ || (flags & 2 /* SHF_ALLOC */)) continue; + put_u64(co, sh + 24, 0); + put_u64(co, sh + 32, kTerabyte); + return co; + } + FAIL("no unallocated PROGBITS section to overstate"); + return co; } // What clang emits for one target: an empty host entry, then the code object. -std::vector valid_bundle(const std::string& target_id, - const std::vector& payload) { - auto b = uncompressed_bundle({{kPayloadOffset, 0, kHostId}, - {kPayloadOffset, payload.size(), target_id}}, +// A non-zero host_size or co_size is declared in that entry instead. HIP never +// loads the host entry, and keeps the size the code object's entry declares. +std::vector target_bundle(const std::string& target_id, + const std::vector& payload, + uint64_t host_size = 0, uint64_t co_size = 0) { + auto b = uncompressed_bundle({{kPayloadOffset, host_size, kHostId}, + {kPayloadOffset, co_size ? co_size : payload.size(), target_id}}, kPayloadOffset + payload.size()); std::memcpy(b.data() + kPayloadOffset, payload.data(), payload.size()); return b; @@ -174,37 +172,30 @@ std::vector marker_payload(const char* marker) { return p; } -// The bundles registered before HIP initialises; the parent rebuilds the same -// bytes to recognise what the capture stored. +// The payloads of the bundles HIP loads no code object from carry a marker, +// which the parent looks for in every file of the archive. Two are registered +// before HIP initialises, one after. constexpr char kPreInitValidMarker[] = "HRR fatbin pre-init valid"; -constexpr char kPreInitLongIdMarker[] = "HRR fatbin pre-init long id"; -constexpr char kPreInitOverflowMarker[] = "HRR fatbin pre-init overflow"; -constexpr uint64_t kPreInitCompressedTag = 0x4852524650524531ull; -constexpr char kPreInitPaddedMarker[] = "HRR fatbin pre-init padded"; -// Takes a bundle past the 1 MiB HIP_HRR_CAPTURE_MAX_BLOB_MB its driver sets. -constexpr size_t kBundlePad = 2u << 20; -// And the ones after, that HIP refuses. +constexpr char kPreInitFarEndMarker[] = "HRR fatbin pre-init far end"; constexpr char kLiveForeignMarker[] = "HRR fatbin live foreign"; -constexpr uint64_t kLiveCompressedTag = 0x485252464C495631ull; +// The four bundles HIP loads, told apart by their kernel. +constexpr char kLiveKernel[] = "hrr_fatbin_live_probe"; +constexpr char kFarEndKernel[] = "hrr_fatbin_far_end_probe"; +constexpr char kOversizedKernel[] = "hrr_fatbin_oversized_probe"; +constexpr char kOverstatedKernel[] = "hrr_fatbin_overstated_probe"; +// Registered before HIP initialises, from a bundle the parent builds: a +// __device__ global, beside a kernel that returns it. +constexpr char kGlobalKernel[] = "hrr_fatbin_global_probe"; +constexpr char kDeviceGlobal[] = "hrr_fatbin_device_global"; +constexpr int kGlobalInit = 0x40000000; +constexpr int kGlobalWritten = 0x3f800000; +constexpr char kGlobalBundleEnv[] = "HRR_FATBIN_GLOBAL_BUNDLE"; std::vector preinit_valid() { return foreign_bundle(kPreInitValidMarker); } -std::vector preinit_compressed() { return short_compressed_bundle(kPreInitCompressedTag); } -std::vector preinit_long_id() { return long_id_bundle(kPreInitLongIdMarker); } -std::vector preinit_overflow() { - return host_overflow_bundle(kForeignId, marker_payload(kPreInitOverflowMarker)); -} -// foreign_bundle() with kBundlePad bytes in its host entry, which HIP never loads. -std::vector preinit_padded() { - const std::vector payload = marker_payload(kPreInitPaddedMarker); - const size_t pad_at = kPayloadOffset + payload.size(); - auto b = uncompressed_bundle({{pad_at, kBundlePad, kHostId}, - {kPayloadOffset, payload.size(), kForeignId}}, - pad_at + kBundlePad); - std::memcpy(b.data() + kPayloadOffset, payload.data(), payload.size()); - return b; +std::vector preinit_far_end() { + return target_bundle(kForeignId, marker_payload(kPreInitFarEndMarker), kTerabyte); } std::vector live_foreign() { return foreign_bundle(kLiveForeignMarker); } -std::vector live_compressed() { return short_compressed_bundle(kLiveCompressedTag); } // Copy a bundle to page-aligned memory that stays mapped for the life of the // process, as a fat binary in .hip_fatbin does: HIP keeps the pointer, and @@ -219,19 +210,47 @@ const void* pin(const std::vector& bytes) { return p; } +// Like pin(), but followed by 64 KiB that cannot be read, so reading the +// terabyte a bundle entry declares faults. +const void* pin_before_guard(const std::vector& bytes) { + constexpr size_t kGuard = 65536; // a multiple of any page size in use + const size_t span = (bytes.size() + kGuard - 1) / kGuard * kGuard; +#if defined(_WIN32) + // Reserved but not committed: no access. + auto* p = static_cast(VirtualAlloc(nullptr, span + kGuard, MEM_RESERVE, + PAGE_NOACCESS)); + REQUIRE(p != nullptr); + REQUIRE(VirtualAlloc(p, span, MEM_COMMIT, PAGE_READWRITE) == p); +#else + void* m = mmap(nullptr, span + kGuard, PROT_READ | PROT_WRITE, + MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); + REQUIRE(m != MAP_FAILED); + auto* p = static_cast(m); + REQUIRE(mprotect(p + span, kGuard, PROT_NONE) == 0); +#endif + std::memcpy(p, bytes.data(), bytes.size()); + return p; +} + void** register_fat_binary(const void* bundle) { // __CudaFatBinaryWrapper; HIP keeps this pointer too. struct Wrapper { uint32_t magic; uint32_t version; const void* binary; const void* dummy; }; return __hipRegisterFatBinary(new Wrapper{kHipfMagic, 1, bundle, nullptr}); } -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; } -)"; +// A kernel writing 1, or with `global`, the value of that __device__ global, +// compiled with `options`. +std::vector compile_code_object(const std::string& kernel, + const std::string& global = "", + std::vector options = {}) { + const std::string src = + (global.empty() ? std::string() : "extern \"C\" __device__ int " + global + " = " + + std::to_string(kGlobalInit) + ";\n") + + "extern \"C\" __global__ void " + kernel + "(int* out) { out[0] = " + + (global.empty() ? std::string("1") : global) + "; }\n"; hiprtcProgram prog = nullptr; - HRR_HIPRTC_CHECK(hiprtcCreateProgram(&prog, src, "hrr_fatbin_probe.hip", 0, nullptr, - nullptr)); + HRR_HIPRTC_CHECK(hiprtcCreateProgram(&prog, src.c_str(), "hrr_fatbin_probe.hip", 0, + nullptr, nullptr)); if (hiprtcCompileProgram(prog, static_cast(options.size()), options.data()) != HIPRTC_SUCCESS) { size_t log_sz = 0; @@ -249,10 +268,23 @@ extern "C" __global__ void hrr_fatbin_probe(int* out) { out[0] = 1; } return co; } -uint64_t read_u64(const std::vector& b, size_t at) { - uint64_t v = 0; - std::memcpy(&v, b.data() + at, sizeof(v)); - return v; +// What clang's module constructor does for a kernel: only the stub's address matters. +void register_function(void** handle, const void* stub, const char* kernel) { + __hipRegisterFunction(handle, stub, const_cast(kernel), kernel, + static_cast(-1), nullptr, nullptr, nullptr, nullptr, nullptr); +} + +// Run the kernel registered for `stub` and return what it wrote. +int launch_probe(const void* stub) { + int* out = nullptr; + HRR_HIP_CHECK(hipMalloc(&out, sizeof(int))); + HRR_HIP_CHECK(hipMemset(out, 0, sizeof(int))); + void* args[] = {&out}; + HRR_HIP_CHECK(hipLaunchKernel(stub, dim3(1), dim3(1), args, 0, nullptr)); + int host = 0; + HRR_HIP_CHECK(hipMemcpy(&host, out, sizeof(int), hipMemcpyDeviceToHost)); + HRR_HIP_CHECK(hipFree(out)); + return host; } } // namespace @@ -264,49 +296,49 @@ TEST_CASE("Unit_HRR_FatBinHeader_Direct", "[.][hrr-direct]") { // HIP is not initialised yet: it only notes these and returns a handle, and // capture sees them when it starts, by sweeping HIP's list. REQUIRE(register_fat_binary(pin(preinit_valid())) != nullptr); - REQUIRE(register_fat_binary(pin(preinit_compressed())) != nullptr); - REQUIRE(register_fat_binary(pin(preinit_long_id())) != nullptr); - REQUIRE(register_fat_binary(pin(preinit_overflow())) != nullptr); - REQUIRE(register_fat_binary(pin(preinit_padded())) != nullptr); + REQUIRE(register_fat_binary(pin(preinit_far_end())) != nullptr); // Initialises HIP, and with it capture. HRR_HIP_CHECK(hipSetDevice(0)); hipDeviceProp_t prop{}; HRR_HIP_CHECK(hipGetDeviceProperties(&prop, 0)); const std::string target = std::string(kTargetPrefix) + prop.gcnArchName; - const std::vector co = compile_code_object(); - REQUIRE(co.size() > 4); - REQUIRE(std::memcmp(co.data(), "\x7f" "ELF", 4) == 0); - - // From here HIP digests each bundle inside the call. These two it loads. - REQUIRE(register_fat_binary(pin(valid_bundle(target, co))) != nullptr); - REQUIRE(register_fat_binary(pin(host_overflow_bundle(target, co))) != nullptr); - // And these it refuses: no code object for this GPU, a bad header. + const std::vector live_co = compile_code_object(kLiveKernel); + const std::vector far_end_co = compile_code_object(kFarEndKernel); + const std::vector oversized_co = compile_code_object(kOversizedKernel); + + // From here HIP digests each bundle inside the call. These four it loads. + void** live = register_fat_binary(pin(target_bundle(target, live_co))); + REQUIRE(live != nullptr); + REQUIRE(register_fat_binary(pin(target_bundle(target, far_end_co, kTerabyte))) != nullptr); + REQUIRE(register_fat_binary(pin_before_guard( + target_bundle(target, oversized_co, 0, kTerabyte))) != nullptr); + REQUIRE(register_fat_binary(pin_before_guard(target_bundle( + target, overstated_section(compile_code_object(kOverstatedKernel)), 0, + kTerabyte))) != nullptr); + // And this one it refuses: no code object for this GPU. REQUIRE(register_fat_binary(pin(live_foreign())) == nullptr); - REQUIRE(register_fat_binary(pin(live_compressed())) == nullptr); + + // A launch for the replay to resolve. + static const char live_stub = 0; + register_function(live, &live_stub, kLiveKernel); + CHECK(launch_probe(&live_stub) == 1); } namespace { -struct RecordedFatBinary { - uint64_t ret = 0; - uint64_t blob_size = 0; - bool has_blob = false; - std::vector blob; -}; +bool contains(const std::vector& hay, const void* needle, size_t n) { + const auto* p = static_cast(needle); + return std::search(hay.begin(), hay.end(), p, p + n) != hay.end(); +} -// True when the capture stored bytes from the start of `bundle`: all of it, or -// the leading part a header overstating or understating its size would give. -bool stored_from(const RecordedFatBinary& r, const std::vector& bundle) { - return r.has_blob && !r.blob.empty() && r.blob.size() <= bundle.size() && - std::memcmp(r.blob.data(), bundle.data(), r.blob.size()) == 0; +bool contains(const std::vector& hay, const char* needle) { + return contains(hay, needle, std::strlen(needle)); } } // namespace // --------------------------------------------------------------------------- -// One capture covers both paths. Before the header checks, every altered -// bundle here was stored: the short compressed headers as their first 23 -// bytes, the others whole, and the two HIP refused after it initialised as -// well, although no launch from them can succeed. +// One capture covers both paths, and its archive replays, also with the live +// bundle stored as a blob. // --------------------------------------------------------------------------- HRR_TEST_CASE(Unit_HRR_FatBinHeaderChecks) { ScopedDir cap(fs::temp_directory_path() / "hrr_fatbin_header.hrr"); @@ -314,9 +346,6 @@ HRR_TEST_CASE(Unit_HRR_FatBinHeaderChecks) { { hrr::test::SpawnProc proc(hrr_self_exe(), /*capture_stdout=*/true, /*capture_stderr=*/true); proc.setEnv("HIP_HRR_CAPTURE_OUTPUT", cap.path.string()); - // The refusals are reported as warnings. - proc.setEnv("AMD_LOG_LEVEL", "2"); - proc.setEnv("HIP_HRR_CAPTURE_MAX_BLOB_MB", "1"); set_proc_search_path(proc); int ret = proc.run("\"Unit_HRR_FatBinHeader_Direct\""); out = proc.getOutput(); @@ -324,123 +353,216 @@ HRR_TEST_CASE(Unit_HRR_FatBinHeaderChecks) { REQUIRE(ret == 0); } INFO("Capture output:\n" << out); + const fs::path archive = hrr_single_process_archive(cap.path); hrr::Archive arc; - REQUIRE(hrr::load_archive(hrr_single_process_archive(cap.path).string(), arc)); + REQUIRE(hrr::load_archive(archive.string(), arc)); - std::vector recorded; + // Every registration is recorded, none with a blob: the three the workload + // made with a null handle (two before HIP initialised, one HIP refused), + // the four HIP loaded with a real one, each naming its code object but the + // one whose ELF overstates a section. + size_t null_handle = 0, loaded = 0; for (const auto& ev : arc.events) { if (ev.header().event_type != HRR_API_HIPREGISTERFATBINARY) continue; REQUIRE(ev.raw_payload.size() >= sizeof(hrr_args___hipRegisterFatBinary)); const auto* a = reinterpret_cast(ev.raw_payload.data()); - RecordedFatBinary r; - r.ret = a->ret; - r.blob_size = a->blob_size; - r.has_blob = a->blob_hash_lo != 0 || a->blob_hash_hi != 0; - if (r.has_blob) REQUIRE(hrr::read_blob(arc, a->blob_hash_lo, a->blob_hash_hi, r.blob)); - recorded.push_back(std::move(r)); + CHECK(a->blob_size == 0); + if (a->blob_hash_lo || a->blob_hash_hi) + CHECK(arc.code_objects.count(hrr::hash_hex(a->blob_hash_lo, a->blob_hash_hi)) == 1); + ++(a->ret ? loaded : null_handle); + } + CHECK(loaded == 4); + CHECK(null_handle >= 3); + + std::vector> stored; + for (const auto* files : {&arc.blobs, &arc.code_objects}) { + for (const auto& [hex, path] : *files) { + std::vector bytes; + REQUIRE(hrr::read_file_capped(path, bytes)); + stored.push_back(std::move(bytes)); + } } - // Registered before HIP initialised (handle recorded as 0): a bundle with a - // sound header is stored whole, even with no code object for this GPU. - const auto pre_valid = preinit_valid(); - size_t pre_valid_hits = 0; - for (const auto& r : recorded) { - if (r.ret == 0 && r.has_blob && r.blob == pre_valid) { - CHECK(r.blob_size == pre_valid.size()); - ++pre_valid_hits; + // The code object of each bundle HIP loaded is there, an ELF on its own + // rather than inside a bundle, and only as long as the ELF: it ends with its + // section header table, as the linker lays it out, even where the bundle + // entry declared a terabyte. A guard that cannot be read follows the bundle, so + // a capture that hashed the declared size did not get this far, nor one that + // trusted the ELF's overstated section, whose code object is not stored. + for (const char* kernel : {kLiveKernel, kFarEndKernel, kOversizedKernel}) { + INFO("Kernel: " << kernel); + size_t hits = 0; + for (const auto& bytes : stored) { + if (!contains(bytes, kernel)) continue; + ++hits; + REQUIRE(bytes.size() >= 64); + CHECK(std::memcmp(bytes.data(), "\x7f" "ELF", 4) == 0); + uint64_t shoff = 0; + uint16_t shentsize = 0, shnum = 0; + std::memcpy(&shoff, bytes.data() + 0x28, sizeof(shoff)); + std::memcpy(&shentsize, bytes.data() + 0x3a, sizeof(shentsize)); + std::memcpy(&shnum, bytes.data() + 0x3c, sizeof(shnum)); + CHECK(shoff + uint64_t(shentsize) * shnum == bytes.size()); } + CHECK(hits == 1); } - CHECK(pre_valid_hits == 1); - - // A bundle over HIP_HRR_CAPTURE_MAX_BLOB_MB is stored whole too, and leaves - // the archive complete. - const auto pre_padded = preinit_padded(); - size_t pre_padded_hits = 0; - for (const auto& r : recorded) - if (r.ret == 0 && r.has_blob && r.blob == pre_padded) ++pre_padded_hits; - CHECK(pre_padded_hits == 1); - CHECK(arc.complete); - - // Registered after HIP initialised and loaded by it (a real handle): the - // valid bundle is stored whole, the one with the wrapping host entry not at - // all, though HIP loaded it, because no end can be computed for it. - std::vector loaded_live; - for (const auto& r : recorded) - if (r.ret != 0) loaded_live.push_back(&r); - REQUIRE(loaded_live.size() == 2); - const RecordedFatBinary* valid = nullptr; - const RecordedFatBinary* overflow = nullptr; - for (const auto* r : loaded_live) (r->has_blob ? valid : overflow) = r; - REQUIRE(valid != nullptr); - REQUIRE(overflow != nullptr); - CHECK(overflow->blob_size == 0); - - // The stored valid bundle is the one the workload built: its code object - // entry ends exactly where the blob ends. - const auto& vb = valid->blob; - REQUIRE(vb.size() > kPayloadOffset + 4); - CHECK(std::memcmp(vb.data(), kBundleMagic, kBundleMagicLen) == 0); - CHECK(read_u64(vb, kBundleMagicLen) == 2); - const size_t second = kTableStart + kEntryHeaderSize + (sizeof(kHostId) - 1); - CHECK(read_u64(vb, second) == kPayloadOffset); - CHECK(read_u64(vb, second + 8) == vb.size() - kPayloadOffset); - CHECK(valid->blob_size == vb.size()); - CHECK(std::memcmp(vb.data() + kPayloadOffset, "\x7f" "ELF", 4) == 0); - - // Refused by HIP after it initialised: the register event is still written, - // without a blob, so replay treats it as a no-op. - size_t refused_live = 0; - for (const auto& r : recorded) - if (r.ret == 0 && !r.has_blob && r.blob_size == 0) ++refused_live; - CHECK(refused_live == 2); - - // No bytes from any altered or refused bundle reached the archive. - const struct { - const char* name; - std::vector bytes; - } rejected[] = { - {"pre-init compressed totalSize below its header", preinit_compressed()}, - {"pre-init entry table past 4096 bytes", preinit_long_id()}, - {"pre-init host entry whose offset plus size wraps", preinit_overflow()}, - {"live bundle with no code object for this GPU", live_foreign()}, - {"live compressed totalSize below its header", live_compressed()}, + + // No file in the archive holds a bundle, or bytes of one HIP never loaded. + for (const auto& bytes : stored) { + CHECK_FALSE(contains(bytes, kBundleMagic, kBundleMagicLen)); + for (const char* marker : {kPreInitValidMarker, kPreInitFarEndMarker, kLiveForeignMarker}) { + INFO("Marker: " << marker); + CHECK_FALSE(contains(bytes, marker)); + } + CHECK_FALSE(contains(bytes, kOverstatedKernel)); + } + + auto [rc, replay] = hrr_playback_merged(archive); + INFO("Replay:\n" << replay); + CHECK(rc == 0); + + // A registration whose bundle is stored as a blob replays through the + // register handler, which loads the blob as a module. Rewrite the live + // kernel's registration that way, and its launch to name no code object, so + // only that module holds the kernel. + std::string events = read_text_file(archive / "events.bin"); + auto at = [&](const hrr::Event& ev) { + const size_t pos = events.find(std::string(ev.raw_payload.begin(), ev.raw_payload.end())); + REQUIRE(pos != std::string::npos); + return pos; }; - for (const auto& bad : rejected) { - INFO("Bundle: " << bad.name); - size_t hits = 0; - for (const auto& r : recorded) - if (stored_from(r, bad.bytes)) ++hits; - CHECK(hits == 0); + std::string hex; + for (const auto& ev : arc.events) { + if (!ev.kernel_launch || ev.kernel_launch->kernel_name != kLiveKernel) continue; + hex = hrr::hash_hex(ev.kernel_launch->co_hash_lo, ev.kernel_launch->co_hash_hi); + // After the stream, the name's length and the name. + const uint64_t none[2] = {0, 0}; + std::memcpy(&events[at(ev) + sizeof(hrr_event_header) + 10 + std::strlen(kLiveKernel)], + none, sizeof(none)); + } + INFO("Code object: " << hex); + REQUIRE(arc.code_objects.count(hex) == 1); + std::vector co; + REQUIRE(hrr::read_file_capped(arc.code_objects.at(hex), co)); + hipDeviceProp_t prop{}; + HRR_HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + const auto bundle = target_bundle(std::string(kTargetPrefix) + prop.gcnArchName, co); + const fs::path blob = archive / "blobs" / hex.substr(0, 2) / (hex + ".blob"); + fs::create_directories(blob.parent_path()); + std::ofstream(blob, std::ios::binary) + .write(reinterpret_cast(bundle.data()), + static_cast(bundle.size())); + fs::remove(arc.code_objects.at(hex)); + const uint64_t blob_size = bundle.size(); + size_t registrations = 0; + 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 (hrr::hash_hex(a->blob_hash_lo, a->blob_hash_hi) != hex) continue; + std::memcpy(&events[at(ev) + offsetof(hrr_args___hipRegisterFatBinary, blob_size)], + &blob_size, sizeof(blob_size)); + ++registrations; } + REQUIRE(registrations == 1); + std::ofstream(archive / "events.bin", std::ios::binary | std::ios::trunc) + .write(events.data(), static_cast(events.size())); + auto [rc_old, replay_old] = hrr_playback_merged(archive); + INFO("Replay of the rewritten archive:\n" << replay_old); + CHECK(rc_old == 0); +} + +// =========================================================================== +// A __device__ global registered before HIP initialises, copied to with +// nothing launched. Capture records its code object only if it resolves +// the global before it sweeps the fat binaries, since resolving it is what +// makes HIP load that code object. +// =========================================================================== +namespace { +// Register a __device__ global from the bundle at kGlobalBundleEnv's path, +// before HIP initialises so that capture records it; the parent compiled it, +// since compiling initialises HIP. +void register_global_bundle(int* shadow) { + const char* path = std::getenv(kGlobalBundleEnv); + if (!path) HRR_SKIP(kGlobalBundleEnv << " is not set"); + std::vector bytes; + REQUIRE(hrr::read_file_capped(path, bytes)); + void** handle = register_fat_binary(pin(bytes)); + REQUIRE(handle != nullptr); + __hipRegisterVar(handle, shadow, const_cast(kDeviceGlobal), + const_cast(kDeviceGlobal), 0, sizeof(int), 0, 0); +} + +// The bundle register_global_bundle() reads, for device 0's GPU. +void write_global_bundle(const fs::path& path) { + hipDeviceProp_t prop{}; + HRR_HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + const auto b = target_bundle(std::string(kTargetPrefix) + prop.gcnArchName, + compile_code_object(kGlobalKernel, kDeviceGlobal)); + std::ofstream(path, std::ios::binary) + .write(reinterpret_cast(b.data()), static_cast(b.size())); +} +} // namespace - // Each refusal says why. - CHECK(out.find("is smaller than its 24-byte header") != std::string::npos); - CHECK(out.find("past the first 4096 bytes") != std::string::npos); - CHECK(out.find("an offset and size that overflow") != std::string::npos); - CHECK(out.find("HIP did not register it") != std::string::npos); +TEST_CASE("Unit_HRR_FatBinGlobalCopy_Direct", "[.][hrr-direct]") { + static int shadow = 0; + register_global_bundle(&shadow); + HRR_HIP_CHECK(hipSetDevice(0)); + HRR_HIP_CHECK(hipMemcpyToSymbol(shadow, &kGlobalWritten, sizeof(int))); +} + +HRR_TEST_CASE(Unit_HRR_FatBinGlobalCopy) { + ScopedDir cap(fs::temp_directory_path() / "hrr_fatbin_global_copy.hrr"); + ScopedDir work(fs::temp_directory_path() / "hrr_fatbin_global_copy.work"); + fs::create_directories(work.path); + const fs::path global_bundle = work.path / "global.bundle"; + write_global_bundle(global_bundle); + hrr_capture_direct("Unit_HRR_FatBinGlobalCopy_Direct", cap.path, 1, + {{kGlobalBundleEnv, global_bundle.string()}}); + const fs::path archive = hrr_single_process_archive(cap.path); + hrr::Archive arc; + REQUIRE(hrr::load_archive(archive.string(), arc)); + // The global's code object is in the archive. + size_t holding = 0; + for (const auto& [hex, path] : arc.code_objects) { + std::vector bytes; + REQUIRE(hrr::read_file_capped(path, bytes)); + holding += contains(bytes, kDeviceGlobal); + } + CHECK(holding == 1); + auto [rc, out] = hrr_playback_merged(archive); + INFO("Replay:\n" << out); + CHECK(rc == 0); } // =========================================================================== // Code objects are stored without their build paths and compiler build. // =========================================================================== +namespace { +// One probe is registered in a fat binary, the other loaded as a module, so +// each leaves its own code object. +constexpr char kScrubFatBinKernel[] = "hrr_scrub_fatbin_probe"; +constexpr char kScrubModuleKernel[] = "hrr_scrub_module_probe"; +} // namespace + 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. + // Both are built with debug information, which is what puts the source file + // name and build directory into a code object. This one is registered in a + // fat binary, of which capture stores the code object HIP loads. + REQUIRE(register_fat_binary(pin(target_bundle( + target, compile_code_object(kScrubFatBinKernel, "", {"-g"})))) != nullptr); + + // Loaded as a module, and launched so replay runs the stored copy. + const std::vector co = compile_code_object(kScrubModuleKernel, "", {"-g"}); hipModule_t mod = nullptr; HRR_HIP_CHECK(hipModuleLoadData(&mod, co.data())); hipFunction_t fn = nullptr; - HRR_HIP_CHECK(hipModuleGetFunction(&fn, mod, "hrr_fatbin_probe")); + HRR_HIP_CHECK(hipModuleGetFunction(&fn, mod, kScrubModuleKernel)); int* d = nullptr; HRR_HIP_CHECK(hipMalloc(&d, sizeof(int))); HRR_HIP_CHECK(hipMemset(d, 0, sizeof(int))); @@ -495,10 +617,6 @@ bool count_scrub_sections(const uint8_t* p, size_t size, ScrubCount& c) { } return true; } - -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) { @@ -507,44 +625,39 @@ HRR_TEST_CASE(Unit_HRR_CodeObjectScrub) { 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; - REQUIRE(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); + REQUIRE(hrr::load_archive(hrr_single_process_archive(cap.path).string(), arc)); - // The bundle's entries are walked from its own table, which zeroing in place - // left as it was. - ScrubCount bundle_count; + // The code objects the loaded fat binaries' registrations name. + std::vector registered; 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) - REQUIRE(count_scrub_sections(blob.data() + off, len, bundle_count)); + if (a->ret && (a->blob_hash_lo || a->blob_hash_hi)) + registered.push_back(hrr::hash_hex(a->blob_hash_lo, a->blob_hash_hi)); + } + + // Every stored code object is scrubbed, the module's and the one HIP took + // from the fat binary alike; each probe is counted where it is found. + ScrubCount module_count, bundle_count; + size_t module_files = 0, bundle_files = 0; + for (const auto& [hex, path] : arc.code_objects) { + INFO("Code object " << path); + std::vector image; + REQUIRE(hrr::read_file_capped(path, image)); + CHECK_FALSE(contains(image, "clang version")); + CHECK_FALSE(contains(image, "hrr_fatbin_probe.hip")); + if (contains(image, kScrubFatBinKernel)) { + CHECK(std::find(registered.begin(), registered.end(), hex) != registered.end()); + ++bundle_files; + REQUIRE(count_scrub_sections(image.data(), image.size(), bundle_count)); + } else { + if (contains(image, kScrubModuleKernel)) ++module_files; + REQUIRE(count_scrub_sections(image.data(), image.size(), module_count)); } - CHECK_FALSE(contains(blob, "clang version")); - CHECK_FALSE(contains(blob, "hrr_fatbin_probe.hip")); } + CHECK(module_files == 1); + CHECK(bundle_files == 1); INFO("Module code objects: " << module_count.elves << " ELF, " << module_count.sections << " sections, " << module_count.left << " not zeroed, " diff --git a/projects/hrr/tools/gen_hrr_api_args.py b/projects/hrr/tools/gen_hrr_api_args.py index 056f5cfd230f..93689dcaea0f 100644 --- a/projects/hrr/tools/gen_hrr_api_args.py +++ b/projects/hrr/tools/gen_hrr_api_args.py @@ -942,7 +942,7 @@ def _windows_exported_symbols() -> Set[str]: # Maps API name -> list of (field_type, field_name, comment) to append EXTRA_FIELDS: Dict[str, List[Tuple[str, str, str]]] = { "__hipRegisterFatBinary": [ - ("uint64_t", "blob_hash_lo", "FNV-1a-128 lo of fat binary"), + ("uint64_t", "blob_hash_lo", "FNV-1a-128 lo of fat binary, or of its code object if blob_size is 0"), ("uint64_t", "blob_hash_hi", "FNV-1a-128 hi"), ("uint64_t", "blob_size", "fat binary byte count"), ],