From 4d319f32a5421e206708478725d8d4630503752e Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Sun, 22 Mar 2026 23:26:40 -0500 Subject: [PATCH 01/12] Add lazy loading support for addFile MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit When addFile is called with lazy=true, only kernel names are indexed from the ELF symbol table (.symtab) — no disassembly runs. Disassembly and source mapping happen on demand when a kernel is first accessed via getKernel(), getKernelLines(), getInstructionsForLine(), or getKernelArguments(). Changes: - kernelDB::addFile: new lazy parameter; indexes via getKernelNamesFromElf - kernelDB::getKernelNamesFromElf: reads .symtab for STT_FUNC/STT_AMDGPU_HSA_KERNEL symbols - kernelDB::ensureKernelLoaded: on-demand disassembly + source mapping for lazy kernels - getKernels() includes lazy kernel names; hasKernel() checks both maps - Python API: KernelDB(lazy=True), add_file(), scan_code_object(), has_kernel() - pybind11: expose addFile, scanCodeObject, hasKernel - pyproject.toml: fix license field format - Add examples/07_lazy_load/ Co-Authored-By: Claude Opus 4.6 --- examples/07_lazy_load/example.py | 99 ++++++++ .../example_runtime_error_failures.py | 40 ++++ include/kernelDB.h | 8 +- kerneldb/api.py | 59 ++++- pyproject.toml | 2 +- src/kernelDB.cc | 217 ++++++++++++++---- src/pybind11_wrapper.cc | 12 +- 7 files changed, 387 insertions(+), 50 deletions(-) create mode 100644 examples/07_lazy_load/example.py create mode 100644 examples/07_lazy_load/example_runtime_error_failures.py diff --git a/examples/07_lazy_load/example.py b/examples/07_lazy_load/example.py new file mode 100644 index 0000000..42062d1 --- /dev/null +++ b/examples/07_lazy_load/example.py @@ -0,0 +1,99 @@ +#!/usr/bin/env python3 +"""Lazy-load example: analyze torch's HIP library with KernelDB (lazy=True).""" + +import sys +import time +from pathlib import Path + +import torch +from kerneldb import KernelDB + + +def get_torch_hip_lib(): + lib_dir = Path(torch.__file__).resolve().parent / "lib" + for name in ("libtorch_hip.so", "libtorch.so"): + p = lib_dir / name + if p.exists(): + return str(p) + raise FileNotFoundError(f"No torch HIP lib in {lib_dir}") + + +def main(): + print("Lazy-load KernelDB Example (torch lib)") + print("=" * 80) + + torch_lib = get_torch_hip_lib() + print(f"Torch lib: {torch_lib}") + + print("\nAnalyzing with KernelDB (lazy=True)...") + kdb = KernelDB(lazy=True) + + t0 = time.perf_counter() + ok = kdb.add_file(torch_lib) + t_add = time.perf_counter() - t0 + if not ok: + print("add_file failed") + return 1 + print(f" add_file (index only): {t_add:.2f} s") + + kernels = kdb.get_kernels() + print(f" get_kernels(): {len(kernels)} names") + print(f"Found {len(kernels)} kernel(s) (showing first 3)") + + shown = 0 + for kernel_name in kernels: + if shown >= 3: + break + try: + kernel = kdb.get_kernel(kernel_name) + except RuntimeError: + continue # ELF name may not match disassembly name + shown += 1 + print(f"\n{'='*80}") + print(f"Kernel: {kernel_name}") + print("=" * 80) + + asm_lines = [] + for block in kernel.get_basic_blocks(): + for inst in block.get_instructions(): + if inst.disassembly: + asm_lines.append(inst.disassembly) + max_asm = 25 + print(f"\nDisassembly (first {min(max_asm, len(asm_lines))} of {len(asm_lines)} lines):") + for ln in asm_lines[:max_asm]: + print(f" {ln}") + if len(asm_lines) > max_asm: + print(f" ... ({len(asm_lines) - max_asm} more lines)") + + lines = kdb.get_kernel_lines(kernel_name) + if not lines: + print("\n(no source line mapping)") + continue + + print(f"\nSource lines: {len(lines)} (range: {min(lines)}-{max(lines)})") + print(f"Basic blocks: {len(kernel.get_basic_blocks())}") + + print("\nInstructions by source line:") + for line in lines: + instructions = kdb.get_instructions_for_line(kernel_name, line) + if instructions: + print(f"\n Line {line}: {len(instructions)} instruction(s)") + for inst in instructions: + print(f" [{inst.line}:{inst.column}] {inst.disassembly}") + + print("\nMemory operations (load/store):") + mem_count = 0 + for line in lines: + mem_ops = kdb.get_instructions_for_line(kernel_name, line, ".*(load|store).*") + mem_count += len(mem_ops) + for inst in mem_ops: + print(f" {inst.disassembly}") + print(f"Total: {mem_count} memory operations") + + print(f"\n{'='*80}") + print("Analysis complete!") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/examples/07_lazy_load/example_runtime_error_failures.py b/examples/07_lazy_load/example_runtime_error_failures.py new file mode 100644 index 0000000..82ab34a --- /dev/null +++ b/examples/07_lazy_load/example_runtime_error_failures.py @@ -0,0 +1,40 @@ +#!/usr/bin/env python3 +"""Show concrete examples of RuntimeError when get_kernel(elf_name) fails (ELF vs disassembly name mismatch).""" + +from pathlib import Path +import torch +from kerneldb import KernelDB + + +def get_torch_hip_lib(): + lib_dir = Path(torch.__file__).resolve().parent / "lib" + for name in ("libtorch_hip.so", "libtorch.so"): + p = lib_dir / name + if p.exists(): + return str(p) + raise FileNotFoundError(f"No torch HIP lib in {lib_dir}") + + +def main(): + kdb = KernelDB(lazy=True) + kdb.add_file(get_torch_hip_lib()) + kernels = kdb.get_kernels() + + max_check = 50 + max_show = 5 + print(f"Checking first {max_check} kernel names for RuntimeError...\n") + failures = [] + for name in list(kernels)[:max_check]: + try: + kdb.get_kernel(name) + except RuntimeError as e: + failures.append((name, str(e))) + print(f"Failures: {len(failures)} / {max_check}\n") + for i, (name, err) in enumerate(failures[:max_show]): + print(f"--- Failure {i + 1} ---") + print(f"ELF name: {name}") + print(f"Error: {err}\n") + + +if __name__ == "__main__": + main() diff --git a/include/kernelDB.h b/include/kernelDB.h index 3700833..a80e970 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -210,7 +210,7 @@ class __attribute__((visibility("default"))) kernelDB { ~kernelDB(); bool getBasicBlocks(const std::string& name, std::vector&); CDNAKernel& getKernel(const std::string& name); - bool addFile(const std::string& name, hsa_agent_t agent, const std::string& strFilter); + bool addFile(const std::string& name, hsa_agent_t agent, const std::string& strFilter, bool lazy = false); bool parseDisassembly(const std::string& text); void mapDisassemblyToSource(hsa_agent_t agent, const char *elfFilePath); bool addKernel(std::unique_ptr kernel); @@ -227,6 +227,10 @@ class __attribute__((visibility("default"))) kernelDB { bool scanCodeObject(const std::string& co_file); bool hasKernel(const std::string& name); private: + /// Get kernel symbol names from a .hsaco ELF without disassembling (reads .symtab). + static std::vector getKernelNamesFromElf(const std::string& fileName); + /// If kernel is lazy-loaded, disassemble its code object and fill kernels_; then remove from lazy set. + void ensureKernelLoaded(const std::string& name); void buildLineMap(size_t offset, size_t hsaco_length, const char *elfFilePath); void extractArgumentsFromDwarf(hsa_agent_t agent, const char *elfFilePath, bool resolve_typedefs); void processKernelsWithAddressMap(const std::map& addrMap); @@ -240,6 +244,8 @@ class __attribute__((visibility("default"))) kernelDB { std::string fileName_; std::map> file_map_; std::set scanned_code_objects_; + /// Lazy-loaded kernels: name -> (hsaco_path, logical_file_name). Filled by addFile(..., lazy=true). + std::map> lazy_kernels_; std::shared_mutex mutex_; }; diff --git a/kerneldb/api.py b/kerneldb/api.py index 29b9503..bc1293e 100644 --- a/kerneldb/api.py +++ b/kerneldb/api.py @@ -38,14 +38,16 @@ class KernelDB: to query instruction-level information mapped to source lines. """ - def __init__(self, binary_path: Optional[str] = None, agent_id: Optional[int] = None): + def __init__(self, binary_path: Optional[str] = None, agent_id: Optional[int] = None, lazy: bool = False): """ Initialize KernelDB Args: - binary_path: Path to HSACO or HIP binary file. If empty string or None, - will search in the running process for fat binaries. + binary_path: Path to HSACO or HIP binary file. If None and lazy=False, + empty string loads the running process executable + shared libs. agent_id: HSA agent handle (if None, will use first GPU) + lazy: If True, create an empty DB; load binaries later with add_file() to avoid + loading the whole process. If False, load binary_path (or process) in constructor. """ # Initialize HSA status = _kerneldb.hsa_init() @@ -61,10 +63,53 @@ def __init__(self, binary_path: Optional[str] = None, agent_id: Optional[int] = if self.agent.handle == 0: raise RuntimeError("No GPU agent found") - # Create kernelDB instance (analysis happens in constructor) - binary_path = binary_path or "" - self._kdb = _kerneldb.KernelDB(self.agent, binary_path) - self.binary_path = binary_path + if lazy: + # Empty DB; add files with add_file() for lazy loading + self._kdb = _kerneldb.KernelDB(self.agent) + self.binary_path = None + else: + binary_path = binary_path or "" + self._kdb = _kerneldb.KernelDB(self.agent, binary_path) + self.binary_path = binary_path or None + + def add_file(self, path: str, filter: str = "", lazy: bool = True) -> bool: + """ + Add a binary (HIP executable or .hsaco). + + With lazy=True (default): only indexes kernel names and their code-object + locations—no disassembly. Disassembly is done on demand when you call + get_kernel(), get_kernel_lines(), get_instructions_for_line(), or access + assembly/arguments for a kernel. + + With lazy=False: full load (disassemble all code objects and map to source), + same as the previous behavior. + + Args: + path: Path to HIP fat binary or .hsaco file + filter: Optional kernel name filter (currently unused in C++) + lazy: If True (default), only index; disassemble on first use per code object. + + Returns: + True on success + """ + return self._kdb.add_file(path, self.agent, filter, lazy) + + def scan_code_object(self, co_file: str) -> bool: + """ + Scan a single .hsaco code object (disassembly + DWARF + args). + Idempotent if the code object was already scanned. + + Args: + co_file: Path to a .hsaco file + + Returns: + True on success + """ + return self._kdb.scan_code_object(co_file) + + def has_kernel(self, name: str) -> bool: + """Return True if a kernel with the given name exists.""" + return self._kdb.has_kernel(name) def get_kernels(self) -> List[str]: """ diff --git a/pyproject.toml b/pyproject.toml index 88f777b..f02b9d5 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,7 +12,7 @@ description = "C++ library for querying data within CDNA kernel implementations, authors = [ { name = "Keith Lowery", email = "Keith.Lowery@amd.com" }, ] -license-files = ["LICENSE"] +license = { text = "MIT" } readme = "README.md" requires-python = ">=3.8" dependencies = [] diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 731f43c..9a3e85d 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -256,14 +256,12 @@ kernelDB::~kernelDB() CDNAKernel& kernelDB::getKernel(const std::string& name) { + ensureKernelLoaded(name); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(name)); if (it != kernels_.end()) - { return *(it->second.get()); - } - else - throw std::runtime_error(name + " kernel does not exist."); + throw std::runtime_error(name + " kernel does not exist."); } @@ -277,27 +275,59 @@ bool kernelDB::addKernel(std::unique_ptr kernel) bool result = true; std::unique_lock lock(mutex_); std::string strName = kernel.get()->getName(); - // std::cout << "Adding kernel \"" << strName << "\"" << std::endl; - if (kernels_.find(strName) == kernels_.end()) + std::string canonical = getKernelName(strName); // same key form as lazy_kernels_ and getKernel lookup + if (kernels_.find(canonical) == kernels_.end()) { - kernels_[strName] = std::move(kernel); + kernels_[canonical] = std::move(kernel); } else { - // std::cout << "You're adding kernel \"" << strName << "\" which we've seen before. Something may be wrong." << std::endl; - kernels_[strName] = std::move(kernel); + kernels_[canonical] = std::move(kernel); result = false; } return result; } -bool kernelDB::addFile(const std::string& name, hsa_agent_t agent, const std::string& strFilter) +void kernelDB::ensureKernelLoaded(const std::string& name) +{ + std::string hsaco_path, logical_file; + { + std::shared_lock lock(mutex_); + auto it = lazy_kernels_.find(getKernelName(name)); + if (it == lazy_kernels_.end()) + return; + hsaco_path = it->second.first; + logical_file = it->second.second; + } + if (!scanCodeObject(hsaco_path)) + return; + try + { + mapDisassemblyToSource(agent_, logical_file.c_str()); + } + catch (const std::runtime_error& e) + { + std::cerr << "Error mapping source for " << logical_file << ": " << e.what() << std::endl; + } + { + std::unique_lock lock(mutex_); + for (auto it = lazy_kernels_.begin(); it != lazy_kernels_.end(); ) + { + if (it->second.first == hsaco_path) + it = lazy_kernels_.erase(it); + else + ++it; + } + } +} + +bool kernelDB::addFile(const std::string& name, hsa_agent_t agent, const std::string& strFilter, bool lazy) { bool bReturn = true; amd_comgr_data_t executable; bool bValidExecutable = false; std::vector isas = ::kernelDB::getIsaList(agent); - std::cout << "Adding " << name << std::endl; + std::cout << "Adding " << name << (lazy ? " (lazy)" : "") << std::endl; if (name.ends_with(".hsaco")) { @@ -320,29 +350,47 @@ bool kernelDB::addFile(const std::string& name, hsa_agent_t agent, const std::st } } - if(bValidExecutable && isas.size()) + if (!bValidExecutable || !isas.size()) + { + bReturn = false; + return bReturn; + } + + if (lazy) { - // Disassemble and parse each code object to discover all kernels + // Only index: get kernel names from ELF symbol table; no disassembly yet. for (const auto& hsaco : file_map_[name]) { - std::string strDisassembly; - getDisassembly(agent, hsaco, strDisassembly); - parseDisassembly(strDisassembly); + std::vector rawNames = getKernelNamesFromElf(hsaco); + std::unique_lock lock(mutex_); + for (const auto& raw : rawNames) + { + std::string demangled = demangleName(raw.c_str()); + std::string canonical = getKernelName(demangled); + if (canonical.empty()) + continue; + lazy_kernels_[canonical] = {hsaco, name}; + } } + return true; + } - // Then map all kernels to source - try - { - mapDisassemblyToSource(agent, name.c_str()); - } - catch (const std::runtime_error& e) - { - std::cerr << "Error adding " << name << "\n\t" << e.what() << std::endl; - bReturn = false; - } + // Eager: disassemble and parse each code object, then map to source + for (const auto& hsaco : file_map_[name]) + { + std::string strDisassembly; + getDisassembly(agent, hsaco, strDisassembly); + parseDisassembly(strDisassembly); } - else + try + { + mapDisassemblyToSource(agent, name.c_str()); + } + catch (const std::runtime_error& e) + { + std::cerr << "Error adding " << name << "\n\t" << e.what() << std::endl; bReturn = false; + } return bReturn; } @@ -403,7 +451,8 @@ bool kernelDB::scanCodeObject(const std::string& co_file) bool kernelDB::hasKernel(const std::string& name) { std::shared_lock lock(mutex_); - return kernels_.find(getKernelName(name)) != kernels_.end(); + std::string key = getKernelName(name); + return kernels_.find(key) != kernels_.end() || lazy_kernels_.find(key) != lazy_kernels_.end(); } std::string kernelDB::extractKernelName(const std::string& line) @@ -705,6 +754,86 @@ void kernelDB::getElfSectionBits(const std::string &fileName, const std::string throw std::runtime_error("Section not found: " + sectionName); } +// AMDGPU kernel symbol type (STT_LOOS = 10); standard ELF STT_FUNC = 2 +#ifndef STT_AMDGPU_HSA_KERNEL +#define STT_AMDGPU_HSA_KERNEL 10 +#endif + +std::vector kernelDB::getKernelNamesFromElf(const std::string& fileName) +{ + std::vector names; + std::ifstream file(fileName, std::ios::binary); + if (!file) + return names; + + Elf64_Ehdr elfHeader; + file.read(reinterpret_cast(&elfHeader), sizeof(elfHeader)); + if (file.gcount() != sizeof(elfHeader) || memcmp(elfHeader.e_ident, ELFMAG, SELFMAG) != 0) + return names; + + file.seekg(elfHeader.e_shoff, std::ios::beg); + std::vector sectionHeaders(elfHeader.e_shnum); + file.read(reinterpret_cast(sectionHeaders.data()), elfHeader.e_shnum * sizeof(Elf64_Shdr)); + if (file.gcount() != static_cast(elfHeader.e_shnum * sizeof(Elf64_Shdr))) + return names; + + const Elf64_Shdr& shstrtab = sectionHeaders[elfHeader.e_shstrndx]; + std::vector shstrtabData(shstrtab.sh_size); + file.seekg(shstrtab.sh_offset, std::ios::beg); + file.read(shstrtabData.data(), shstrtab.sh_size); + + Elf64_Word textShndx = SHN_UNDEF; + size_t symtabOffset = 0, symtabSize = 0, strtabOffset = 0, strtabSize = 0; + + for (Elf64_Word i = 0; i < elfHeader.e_shnum; ++i) + { + std::string secName(&shstrtabData[sectionHeaders[i].sh_name]); + if (secName == ".text") + textShndx = i; + else if (secName == ".symtab") + { + symtabOffset = sectionHeaders[i].sh_offset; + symtabSize = sectionHeaders[i].sh_size; + } + else if (secName == ".strtab") + { + strtabOffset = sectionHeaders[i].sh_offset; + strtabSize = sectionHeaders[i].sh_size; + } + } + + if (textShndx == SHN_UNDEF || symtabSize == 0 || strtabSize == 0) + return names; + + std::vector strtabData(strtabSize); + file.seekg(strtabOffset, std::ios::beg); + file.read(strtabData.data(), strtabSize); + + const size_t numSyms = symtabSize / sizeof(Elf64_Sym); + file.seekg(symtabOffset, std::ios::beg); + for (size_t s = 0; s < numSyms; ++s) + { + Elf64_Sym sym; + file.read(reinterpret_cast(&sym), sizeof(sym)); + if (file.gcount() != static_cast(sizeof(sym))) + break; + if (sym.st_shndx == SHN_UNDEF) + continue; + uint8_t type = ELF64_ST_TYPE(sym.st_info); + if (type != STT_FUNC && type != STT_AMDGPU_HSA_KERNEL) + continue; + if (sym.st_shndx != textShndx) + continue; + if (sym.st_name >= strtabSize) + continue; + const char* nameStr = &strtabData[sym.st_name]; + if (!nameStr[0]) + continue; + names.push_back(std::string(nameStr)); + } + return names; +} + //using namespace llvm; //using namespace llvm::object; @@ -962,6 +1091,7 @@ void kernelDB::mapDisassemblyToSource(hsa_agent_t agent, const char *elfFilePath std::string kernelDB::getFileName(const std::string& kernel, size_t index) { + ensureKernelLoaded(kernel); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(kernel)); if (it != kernels_.end()) @@ -977,6 +1107,7 @@ std::string kernelDB::getFileName(const std::string& kernel, size_t index) std::vector kernelDB::getInstructionsForLine(const std::string& kernel_name, uint32_t line, const std::string& match) { + ensureKernelLoaded(kernel_name); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(kernel_name)); if (it != kernels_.end()) @@ -987,6 +1118,7 @@ std::vector kernelDB::getInstructionsForLine(const std::string& k const std::vector& kernelDB::getInstructionsForLine(const std::string& kernel_name, uint32_t line) { + ensureKernelLoaded(kernel_name); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(kernel_name)); if (it != kernels_.end()) @@ -998,30 +1130,35 @@ const std::vector& kernelDB::getInstructionsForLine(const std::st void kernelDB::getKernels(std::vector& out) { std::shared_lock lock(mutex_); - auto it = kernels_.begin(); - while (it != kernels_.end()) - { - out.push_back(it->first); - it++; - } + for (const auto& p : kernels_) + out.push_back(p.first); + for (const auto& p : lazy_kernels_) + out.push_back(p.first); } void kernelDB::getKernelLines(const std::string& kernel, std::vector& out) { + ensureKernelLoaded(kernel); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(kernel)); if (it != kernels_.end()) - { - it->second.get()->getLineNumbers(out); - } + it->second.get()->getLineNumbers(out); } std::vector kernelDB::getKernelArguments(const std::string& kernel_name, bool resolve_typedefs) { - // If we need to resolve typedefs and haven't done so yet, re-extract with resolution - if (resolve_typedefs) { - extractArgumentsFromDwarf(agent_, fileName_.c_str(), true); + std::string logical_file; + { + std::shared_lock lock(mutex_); + auto lit = lazy_kernels_.find(getKernelName(kernel_name)); + if (lit != lazy_kernels_.end()) + logical_file = lit->second.second; } + ensureKernelLoaded(kernel_name); + if (resolve_typedefs && !logical_file.empty()) + extractArgumentsFromDwarf(agent_, logical_file.c_str(), true); + else if (resolve_typedefs && !fileName_.empty()) + extractArgumentsFromDwarf(agent_, fileName_.c_str(), true); std::shared_lock lock(mutex_); auto it = kernels_.find(getKernelName(kernel_name)); diff --git a/src/pybind11_wrapper.cc b/src/pybind11_wrapper.cc index fbb2344..072ee86 100644 --- a/src/pybind11_wrapper.cc +++ b/src/pybind11_wrapper.cc @@ -127,8 +127,18 @@ PYBIND11_MODULE(_kerneldb, m) { "Create KernelDB instance with agent and filename", py::arg("agent"), py::arg("filename")) .def(py::init(), - "Create KernelDB instance with agent (searches process)", + "Create empty KernelDB for lazy loading (add files with add_file)", py::arg("agent")) + .def("add_file", &kernelDB::kernelDB::addFile, + "Add a binary. If lazy=True (default), only index kernel names (no disassembly). " + "Disassembly runs on demand when you get assembly/source/args for a kernel.", + py::arg("path"), py::arg("agent"), py::arg("filter") = "", py::arg("lazy") = true) + .def("scan_code_object", &kernelDB::kernelDB::scanCodeObject, + "Scan a single .hsaco code object (disassembly + DWARF + args). Idempotent if already scanned.", + py::arg("co_file")) + .def("has_kernel", &kernelDB::kernelDB::hasKernel, + "Return True if a kernel with the given name exists", + py::arg("name")) .def("get_kernel", &kernelDB::kernelDB::getKernel, "Get a kernel by name", py::arg("name"), From 43a64b789e0f56f6aef12ac313e5363c0f5fb935 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Sun, 22 Mar 2026 22:56:06 -0700 Subject: [PATCH 02/12] Per-kernel lazy loading: only disassemble the requested kernel ensureKernelLoaded() previously called scanCodeObject() which disassembles every kernel in the .hsaco, defeating lazy loading for code objects with many kernels. Now it calls scanCodeObjectForKernel() which uses parseDisassemblyForKernel() to skip non-target kernels, and scopes DWARF/argument processing to the single loaded kernel. The code object is not marked as fully scanned so other kernels from the same .hsaco can still be loaded on demand. Co-Authored-By: Claude Opus 4.6 --- examples/07_lazy_load/example.py | 99 -------- .../example_runtime_error_failures.py | 40 --- include/kernelDB.h | 4 +- src/kernelDB.cc | 240 ++++++++++++++++-- 4 files changed, 225 insertions(+), 158 deletions(-) delete mode 100644 examples/07_lazy_load/example.py delete mode 100644 examples/07_lazy_load/example_runtime_error_failures.py diff --git a/examples/07_lazy_load/example.py b/examples/07_lazy_load/example.py deleted file mode 100644 index 42062d1..0000000 --- a/examples/07_lazy_load/example.py +++ /dev/null @@ -1,99 +0,0 @@ -#!/usr/bin/env python3 -"""Lazy-load example: analyze torch's HIP library with KernelDB (lazy=True).""" - -import sys -import time -from pathlib import Path - -import torch -from kerneldb import KernelDB - - -def get_torch_hip_lib(): - lib_dir = Path(torch.__file__).resolve().parent / "lib" - for name in ("libtorch_hip.so", "libtorch.so"): - p = lib_dir / name - if p.exists(): - return str(p) - raise FileNotFoundError(f"No torch HIP lib in {lib_dir}") - - -def main(): - print("Lazy-load KernelDB Example (torch lib)") - print("=" * 80) - - torch_lib = get_torch_hip_lib() - print(f"Torch lib: {torch_lib}") - - print("\nAnalyzing with KernelDB (lazy=True)...") - kdb = KernelDB(lazy=True) - - t0 = time.perf_counter() - ok = kdb.add_file(torch_lib) - t_add = time.perf_counter() - t0 - if not ok: - print("add_file failed") - return 1 - print(f" add_file (index only): {t_add:.2f} s") - - kernels = kdb.get_kernels() - print(f" get_kernels(): {len(kernels)} names") - print(f"Found {len(kernels)} kernel(s) (showing first 3)") - - shown = 0 - for kernel_name in kernels: - if shown >= 3: - break - try: - kernel = kdb.get_kernel(kernel_name) - except RuntimeError: - continue # ELF name may not match disassembly name - shown += 1 - print(f"\n{'='*80}") - print(f"Kernel: {kernel_name}") - print("=" * 80) - - asm_lines = [] - for block in kernel.get_basic_blocks(): - for inst in block.get_instructions(): - if inst.disassembly: - asm_lines.append(inst.disassembly) - max_asm = 25 - print(f"\nDisassembly (first {min(max_asm, len(asm_lines))} of {len(asm_lines)} lines):") - for ln in asm_lines[:max_asm]: - print(f" {ln}") - if len(asm_lines) > max_asm: - print(f" ... ({len(asm_lines) - max_asm} more lines)") - - lines = kdb.get_kernel_lines(kernel_name) - if not lines: - print("\n(no source line mapping)") - continue - - print(f"\nSource lines: {len(lines)} (range: {min(lines)}-{max(lines)})") - print(f"Basic blocks: {len(kernel.get_basic_blocks())}") - - print("\nInstructions by source line:") - for line in lines: - instructions = kdb.get_instructions_for_line(kernel_name, line) - if instructions: - print(f"\n Line {line}: {len(instructions)} instruction(s)") - for inst in instructions: - print(f" [{inst.line}:{inst.column}] {inst.disassembly}") - - print("\nMemory operations (load/store):") - mem_count = 0 - for line in lines: - mem_ops = kdb.get_instructions_for_line(kernel_name, line, ".*(load|store).*") - mem_count += len(mem_ops) - for inst in mem_ops: - print(f" {inst.disassembly}") - print(f"Total: {mem_count} memory operations") - - print(f"\n{'='*80}") - print("Analysis complete!") - return 0 - - -if __name__ == "__main__": - sys.exit(main()) diff --git a/examples/07_lazy_load/example_runtime_error_failures.py b/examples/07_lazy_load/example_runtime_error_failures.py deleted file mode 100644 index 82ab34a..0000000 --- a/examples/07_lazy_load/example_runtime_error_failures.py +++ /dev/null @@ -1,40 +0,0 @@ -#!/usr/bin/env python3 -"""Show concrete examples of RuntimeError when get_kernel(elf_name) fails (ELF vs disassembly name mismatch).""" - -from pathlib import Path -import torch -from kerneldb import KernelDB - - -def get_torch_hip_lib(): - lib_dir = Path(torch.__file__).resolve().parent / "lib" - for name in ("libtorch_hip.so", "libtorch.so"): - p = lib_dir / name - if p.exists(): - return str(p) - raise FileNotFoundError(f"No torch HIP lib in {lib_dir}") - - -def main(): - kdb = KernelDB(lazy=True) - kdb.add_file(get_torch_hip_lib()) - kernels = kdb.get_kernels() - - max_check = 50 - max_show = 5 - print(f"Checking first {max_check} kernel names for RuntimeError...\n") - failures = [] - for name in list(kernels)[:max_check]: - try: - kdb.get_kernel(name) - except RuntimeError as e: - failures.append((name, str(e))) - print(f"Failures: {len(failures)} / {max_check}\n") - for i, (name, err) in enumerate(failures[:max_show]): - print(f"--- Failure {i + 1} ---") - print(f"ELF name: {name}") - print(f"Error: {err}\n") - - -if __name__ == "__main__": - main() diff --git a/include/kernelDB.h b/include/kernelDB.h index a80e970..d6cc249 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -225,6 +225,8 @@ class __attribute__((visibility("default"))) kernelDB { static void getElfSectionBits(const std::string &fileName, const std::string §ionName, size_t& offset, std::vector& sectionData ); std::vector getKernelArguments(const std::string& kernel_name, bool resolve_typedefs = false); bool scanCodeObject(const std::string& co_file); + bool scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName); + bool parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel); bool hasKernel(const std::string& name); private: /// Get kernel symbol names from a .hsaco ELF without disassembling (reads .symtab). @@ -233,7 +235,7 @@ class __attribute__((visibility("default"))) kernelDB { void ensureKernelLoaded(const std::string& name); void buildLineMap(size_t offset, size_t hsaco_length, const char *elfFilePath); void extractArgumentsFromDwarf(hsa_agent_t agent, const char *elfFilePath, bool resolve_typedefs); - void processKernelsWithAddressMap(const std::map& addrMap); + void processKernelsWithAddressMap(const std::map& addrMap, const std::string& targetKernel = ""); parse_mode getLineType(std::string& line); std::string extractKernelName(const std::string& line); static bool isBranch(const std::string& instruction); diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 9a3e85d..7e75935 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -290,34 +290,25 @@ bool kernelDB::addKernel(std::unique_ptr kernel) void kernelDB::ensureKernelLoaded(const std::string& name) { + std::string canonical = getKernelName(name); std::string hsaco_path, logical_file; { std::shared_lock lock(mutex_); - auto it = lazy_kernels_.find(getKernelName(name)); + auto it = lazy_kernels_.find(canonical); if (it == lazy_kernels_.end()) + { return; + } hsaco_path = it->second.first; logical_file = it->second.second; } - if (!scanCodeObject(hsaco_path)) - return; - try - { - mapDisassemblyToSource(agent_, logical_file.c_str()); - } - catch (const std::runtime_error& e) + if (!scanCodeObjectForKernel(hsaco_path, canonical)) { - std::cerr << "Error mapping source for " << logical_file << ": " << e.what() << std::endl; + return; } { std::unique_lock lock(mutex_); - for (auto it = lazy_kernels_.begin(); it != lazy_kernels_.end(); ) - { - if (it->second.first == hsaco_path) - it = lazy_kernels_.erase(it); - else - ++it; - } + lazy_kernels_.erase(canonical); } } @@ -448,6 +439,60 @@ bool kernelDB::scanCodeObject(const std::string& co_file) return true; } +bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName) +{ + std::string strDisassembly; + if (!getDisassembly(agent_, co_file, strDisassembly)) + { + return false; + } + + parseDisassemblyForKernel(strDisassembly, kernelName); + + std::map addrMap; + buildDwarfAddressMap(co_file.c_str(), 0, 0, addrMap); + + if (!addrMap.empty()) + { + processKernelsWithAddressMap(addrMap, kernelName); + } + + // Extract arguments only for the target kernel + std::map> kernelArgsMap; + extractKernelArguments(co_file.c_str(), 0, 0, kernelArgsMap, false); + + { + std::unique_lock lock(mutex_); + for (const auto& entry : kernelArgsMap) + { + std::string demangledName = demangleName(entry.first.c_str()); + std::string searchName = getKernelName(demangledName); + + auto it = kernels_.find(searchName); + if (it == kernels_.end()) + { + for (auto& k : kernels_) + { + if (k.first.compare(0, searchName.length(), searchName) == 0 && + (k.first.length() == searchName.length() || k.first[searchName.length()] == '(')) + { + k.second.get()->setArguments(entry.second); + break; + } + } + } + else + { + it->second.get()->setArguments(entry.second); + } + } + // Do NOT mark co_file in scanned_code_objects_ — other kernels from this + // code object may still need loading later. + } + + return true; +} + bool kernelDB::hasKernel(const std::string& name) { std::shared_lock lock(mutex_); @@ -709,6 +754,161 @@ bool kernelDB::parseDisassembly(const std::string& text) return bReturn; } +bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel) +{ + bool bReturn = true; + std::istringstream in(text); + std::string line; + parse_mode mode = BEGIN; + std::string strKernel; + uint32_t block_count = 0; + std::unique_ptr kernel; + CDNAKernel *current_kernel = nullptr; + std::unique_ptr block; + basicBlock *current_block = nullptr; + std::map> markers; + getBlockMarkers(text, markers); + std::map>::iterator mit; + bool bDoingKernels = false; + bool skip = false; // true when current kernel is not the target + while(std::getline(in,line)) + { + bool blockCreated = false; + std::vector tokens; + mode = getLineType(line); + switch(mode) + { + case BEGIN: + mode = KERNEL; + block_count = 0; + break; + case KERNEL: + { + bDoingKernels = true; + strKernel = extractKernelName(line); + std::string demangledName = demangleName(strKernel.c_str()); + std::string canonical = getKernelName(demangledName); + if (canonical != targetKernel) + { + skip = true; + current_kernel = nullptr; + current_block = nullptr; + break; + } + skip = false; + kernel = std::make_unique(demangledName); + mit = markers.find(demangledName); + assert(mit != markers.end()); + current_kernel = kernel.get(); + mode = BBLOCK; + addKernel(std::move(kernel)); + break; + } + case BBLOCK: + if (!bDoingKernels || skip) + { + break; + } + split(line, tokens, " ", false); + if (tokens.size()) + { + if (!current_block) + { + block = std::make_unique(); + block_count++; + current_block = block.get(); + blockCreated = true; + } + trim(tokens[0]); + if (isBranch(tokens[0])) + { + if (current_kernel) + { + current_kernel->addBlock(block_count, std::move(block)); + } + else + { + std::cout << "Disassembly parsing error. Processing a branch instruction when there's not a kernel currently defined.\n"; + std::cout << line << std::endl; + abort(); + } + if (tokens[0].find("s_endpgm") != std::string::npos) + { + blockCreated = true; + block = std::make_unique(); + block_count++; + current_block = block.get(); + } + else + { + current_block = nullptr; + continue; + } + } + std::vector inst_tokens; + instruction_t inst; + + if (tokens.size() > 1 && tokens[0].find("_") != std::string::npos) + { + split(tokens[0], inst_tokens, "_", false); + if (inst_tokens.size() > 2 && (inst_tokens[1] == "load" || inst_tokens[1] == "store")) + { + inst.prefix_ = inst_tokens[0]; + inst.type_ = inst_tokens[1]; + inst.size_ = inst_tokens[2]; + } + inst.inst_ = tokens[0]; + inst.disassembly_ = line; + for (size_t i = 1; i < tokens.size() && tokens[i].find("//") == std::string::npos; i++) + { + inst.operands_.push_back(tokens[i]); + } + size_t i = 1; + while (tokens[i].find("//") == std::string::npos) + { + i++; + } + std::string strAddress = tokens[++i]; + strAddress.pop_back(); + inst.address_ = std::stoull(strAddress, nullptr, 16); + if (!blockCreated && mit != markers.end() && (mit->second.find(inst.address_) != mit->second.end())) + { + if (current_block) + { + current_kernel->addBlock(block_count, std::move(block)); + } + block = std::make_unique(); + block_count++; + current_block = block.get(); + } + inst.block_ = current_block; + current_block->addInstruction(inst); + if (inst.inst_ == "s_endpgm") + { + if (current_kernel && current_block) + { + current_kernel->addBlock(block_count, std::move(block)); + current_block = nullptr; + } + else + { + std::cerr << "Error parsing disassembly - s_endpgm without current kernel or block\n"; + } + } + } + } + break; + case BRANCH: + current_block = nullptr; + break; + default: + break; + } + } + + return bReturn; +} + void kernelDB::getElfSectionBits(const std::string &fileName, const std::string §ionName, size_t& offset, std::vector& sectionData ) { std::ifstream file(fileName, std::ios::binary); if (!file) { @@ -987,12 +1187,17 @@ void CDNAKernel::getSourceCode(std::vector& outputLines) } -void kernelDB::processKernelsWithAddressMap(const std::map& addrMap) +void kernelDB::processKernelsWithAddressMap(const std::map& addrMap, const std::string& targetKernel) { std::unique_lock lock(mutex_); auto it = kernels_.begin(); while(it != kernels_.end()) { + if (!targetKernel.empty() && it->first != targetKernel) + { + it++; + continue; + } const auto& blocks = it->second.get()->getBasicBlocks(); for (const auto& block : blocks) { @@ -1018,7 +1223,6 @@ void kernelDB::processKernelsWithAddressMap(const std::mapsecond.get()->addLine(MISSING_SOURCE_INFO, inst); - //std::cout << "No match for " << std::hex << "0x" << instruction.address_ << std::dec << std::endl; } } } From cc748e9b3c018a8a4bd44dabd1953e540771d4ff Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 04:16:30 -0700 Subject: [PATCH 03/12] Use targeted disassembly for lazy kernel loading - Add getDisassemblyForSymbol() using --disassemble-symbols instead of full -d, so scanCodeObjectForKernel() only disassembles the requested kernel - Fix shell quoting in invokeProgram() to handle kernel names with parentheses or other shell metacharacters - Fix trailing space in disassembly_params ("-d " -> "-d") - Accept STT_OBJECT .kd symbols in getKernelNamesFromElf() Co-Authored-By: Claude Opus 4.6 --- include/kernelDB.h | 1 + src/disassemble.cc | 45 +++++++++++++++++++++++++++++++++++++++++---- src/kernelDB.cc | 33 +++++++++++++++++++++++++++++---- 3 files changed, 71 insertions(+), 8 deletions(-) diff --git a/include/kernelDB.h b/include/kernelDB.h index d6cc249..efe9d26 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -113,6 +113,7 @@ struct KernelArgument { bool buildDwarfAddressMap(const char* filename, size_t offset, size_t hsaco_length, std::map& addressMap); SourceLocation getSourceLocation(const std::map& addrMap, Dwarf_Addr addr); __attribute__((visibility("default"))) bool getDisassembly(hsa_agent_t agent, const std::string& fileName, std::string& out); +bool getDisassemblyForSymbol(hsa_agent_t agent, const std::string& fileName, const std::string& symbolName, std::string& out); bool invokeProgram(const std::string& programName, const std::vector& params, const std::string& outputFileName); std::string create_temp_file_segment(const std::string& filename, std::streamoff offset, std::streamsize length); __attribute__((visibility("default"))) std::vector extractCodeObjects(hsa_agent_t agent, const std::string& fileName); diff --git a/src/disassemble.cc b/src/disassemble.cc index 57f5fb6..7d65370 100644 --- a/src/disassemble.cc +++ b/src/disassemble.cc @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include "include/kernelDB.h" -std::vector disassembly_params = {"-d ", "--arch-name=amdgcn"}; +std::vector disassembly_params = {"-d", "--arch-name=amdgcn"}; void readFileToString(const std::string& filename, std::string& content) { std::ifstream file(filename, std::ios::binary); @@ -92,13 +92,50 @@ bool getDisassembly(hsa_agent_t agent, const std::string& fileName, std::string& return true; } +bool getDisassemblyForSymbol(hsa_agent_t agent, const std::string& fileName, + const std::string& symbolName, std::string& out) +{ + // Use --disassemble-symbols to extract only the requested kernel. + std::vector parms = {"--arch-name=amdgcn"}; + char name[64]; + memset(name, 0, sizeof(name)); + hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, name); + if (status != HSA_STATUS_SUCCESS) + return false; + + std::stringstream ss; + ss << "--mcpu=" << name; + parms.push_back(ss.str()); + parms.push_back("--disassemble-symbols=" + symbolName); + parms.push_back(fileName); + + char temp_filename[L_tmpnam]; + if (tmpnam(temp_filename) == nullptr) + throw std::runtime_error("Failed to generate temporary filename"); + + if (invokeProgram(disassembler, parms, temp_filename)) + { + readFileToString(temp_filename, out); + unlink(temp_filename); + return out.length() != 0; + } + return false; +} + bool invokeProgram(const std::string& programName, const std::vector& params, const std::string& outputFileName) { - // Construct the command string + // Construct the command string with shell-safe quoting. + // Parameters are single-quoted so kernel names containing parentheses, + // spaces, or other shell metacharacters are passed through safely. std::stringstream command; command << programName; for (const auto& param : params) { - // Basic escaping of parameters (assumes no spaces in params; enhance if needed) - command << " " << param; + std::string escaped = param; + size_t pos = 0; + while ((pos = escaped.find('\'', pos)) != std::string::npos) { + escaped.replace(pos, 1, "'\\''"); + pos += 4; + } + command << " '" << escaped << "'"; } // Redirect stdout to outputFileName command << " > " << outputFileName; diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 7e75935..fb2648e 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -442,7 +442,26 @@ bool kernelDB::scanCodeObject(const std::string& co_file) bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName) { std::string strDisassembly; - if (!getDisassembly(agent_, co_file, strDisassembly)) + // Use targeted disassembly (--disassemble-symbols) to extract only the + // requested kernel instead of disassembling the entire code object. + // Try the FUNC symbol name first, then with .kd suffix (in case + // the caller passed a kernel descriptor name). + bool gotDisasm = false; + std::string funcName = kernelName; + if (funcName.size() > 3 && funcName.substr(funcName.size() - 3) == ".kd") + funcName = funcName.substr(0, funcName.size() - 3); + for (const auto& sym : {funcName, funcName + ".kd"}) + { + try { + gotDisasm = getDisassemblyForSymbol(agent_, co_file, sym, strDisassembly); + } catch (...) { + gotDisasm = false; + } + if (gotDisasm) + break; + strDisassembly.clear(); + } + if (!gotDisasm) { return false; } @@ -1020,16 +1039,22 @@ std::vector kernelDB::getKernelNamesFromElf(const std::string& file if (sym.st_shndx == SHN_UNDEF) continue; uint8_t type = ELF64_ST_TYPE(sym.st_info); - if (type != STT_FUNC && type != STT_AMDGPU_HSA_KERNEL) + bool isCode = (type == STT_FUNC || type == STT_AMDGPU_HSA_KERNEL); + bool isObject = (type == STT_OBJECT); + if (!isCode && !isObject) continue; - if (sym.st_shndx != textShndx) + if (isCode && sym.st_shndx != textShndx) continue; if (sym.st_name >= strtabSize) continue; const char* nameStr = &strtabData[sym.st_name]; if (!nameStr[0]) continue; - names.push_back(std::string(nameStr)); + std::string symName(nameStr); + // Only keep .kd kernel descriptors from OBJECT symbols + if (isObject && (symName.size() < 3 || symName.substr(symName.size() - 3) != ".kd")) + continue; + names.push_back(symName); } return names; } From 1ade35dec9c1b11616d53b06ffb471dfc317a255 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 04:49:18 -0700 Subject: [PATCH 04/12] Pass raw ELF symbol name to --disassemble-symbols The demangled kernel name can contain commas (e.g. C++ template args) which --disassemble-symbols interprets as multiple symbol names. Store the raw/mangled ELF symbol name in lazy_kernels_ and pass it to scanCodeObjectForKernel/getDisassemblyForSymbol instead. Co-Authored-By: Claude Opus 4.6 --- include/kernelDB.h | 9 +++++++-- src/kernelDB.cc | 12 ++++++------ 2 files changed, 13 insertions(+), 8 deletions(-) diff --git a/include/kernelDB.h b/include/kernelDB.h index efe9d26..cf61a43 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -247,8 +247,13 @@ class __attribute__((visibility("default"))) kernelDB { std::string fileName_; std::map> file_map_; std::set scanned_code_objects_; - /// Lazy-loaded kernels: name -> (hsaco_path, logical_file_name). Filled by addFile(..., lazy=true). - std::map> lazy_kernels_; + struct LazyKernelEntry { + std::string hsaco_path; + std::string logical_file; + std::string elf_symbol; // raw (mangled) ELF symbol name for --disassemble-symbols + }; + /// Lazy-loaded kernels: canonical name -> entry. Filled by addFile(..., lazy=true). + std::map lazy_kernels_; std::shared_mutex mutex_; }; diff --git a/src/kernelDB.cc b/src/kernelDB.cc index fb2648e..f45fee1 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -291,7 +291,7 @@ bool kernelDB::addKernel(std::unique_ptr kernel) void kernelDB::ensureKernelLoaded(const std::string& name) { std::string canonical = getKernelName(name); - std::string hsaco_path, logical_file; + std::string hsaco_path, elf_symbol; { std::shared_lock lock(mutex_); auto it = lazy_kernels_.find(canonical); @@ -299,10 +299,10 @@ void kernelDB::ensureKernelLoaded(const std::string& name) { return; } - hsaco_path = it->second.first; - logical_file = it->second.second; + hsaco_path = it->second.hsaco_path; + elf_symbol = it->second.elf_symbol; } - if (!scanCodeObjectForKernel(hsaco_path, canonical)) + if (!scanCodeObjectForKernel(hsaco_path, elf_symbol)) { return; } @@ -360,7 +360,7 @@ bool kernelDB::addFile(const std::string& name, hsa_agent_t agent, const std::st std::string canonical = getKernelName(demangled); if (canonical.empty()) continue; - lazy_kernels_[canonical] = {hsaco, name}; + lazy_kernels_[canonical] = {hsaco, name, raw}; } } return true; @@ -1381,7 +1381,7 @@ std::vector kernelDB::getKernelArguments(const std::string& kern std::shared_lock lock(mutex_); auto lit = lazy_kernels_.find(getKernelName(kernel_name)); if (lit != lazy_kernels_.end()) - logical_file = lit->second.second; + logical_file = lit->second.logical_file; } ensureKernelLoaded(kernel_name); if (resolve_typedefs && !logical_file.empty()) From bc652ecfb8a1a1b2a1837a7a337b328ce75524cc Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 05:01:54 -0700 Subject: [PATCH 05/12] Fix target kernel name comparison in parseDisassemblyForKernel The targetKernel parameter passed from scanCodeObjectForKernel is the raw (mangled) ELF symbol name, but parseDisassemblyForKernel compares it against demangled canonical names extracted from the disassembly output. This mismatch caused every kernel to be skipped, resulting in 0 parsed instructions. Normalize targetKernel through demangleName + getKernelName at the top of both parseDisassemblyForKernel and processKernelsWithAddressMap so comparisons use the same canonical form on both sides. Co-Authored-By: Claude Opus 4.6 --- src/kernelDB.cc | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/src/kernelDB.cc b/src/kernelDB.cc index f45fee1..af2483d 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -775,6 +775,11 @@ bool kernelDB::parseDisassembly(const std::string& text) bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel) { + // Normalize the target name so it can be compared against demangled canonical + // forms extracted from the disassembly output. The caller may pass a raw + // (mangled) ELF symbol name, so we demangle and canonicalize it here. + std::string targetCanonical = getKernelName(demangleName(targetKernel.c_str())); + bool bReturn = true; std::istringstream in(text); std::string line; @@ -807,7 +812,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str strKernel = extractKernelName(line); std::string demangledName = demangleName(strKernel.c_str()); std::string canonical = getKernelName(demangledName); - if (canonical != targetKernel) + if (canonical != targetCanonical) { skip = true; current_kernel = nullptr; @@ -1214,11 +1219,13 @@ void CDNAKernel::getSourceCode(std::vector& outputLines) void kernelDB::processKernelsWithAddressMap(const std::map& addrMap, const std::string& targetKernel) { + // Normalize the target so raw ELF symbol names match demangled kernel keys. + std::string targetCanonical = targetKernel.empty() ? "" : getKernelName(demangleName(targetKernel.c_str())); std::unique_lock lock(mutex_); auto it = kernels_.begin(); while(it != kernels_.end()) { - if (!targetKernel.empty() && it->first != targetKernel) + if (!targetCanonical.empty() && it->first != targetCanonical) { it++; continue; From e368f6980ec823adf8c9de6a7d1730c446f173e3 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 05:14:54 -0700 Subject: [PATCH 06/12] Fix segfault: guard getLineType against empty lines and EOF getLineType dereferenced line.begin() without checking if the string was empty, causing a segfault when getBlockMarkers looped past EOF on disassembly output that contained headers but no kernel sections. Also fix getBlockMarkers to return early on EOF instead of looping forever, and replace the assert(mit != markers.end()) with a graceful skip so release builds don't crash. Add debug logging to ensureKernelLoaded and scanCodeObjectForKernel for diagnosing lazy loading issues. Co-Authored-By: Claude Opus 4.6 --- src/kernelDB.cc | 28 +++++++++++++++++++++++++--- 1 file changed, 25 insertions(+), 3 deletions(-) diff --git a/src/kernelDB.cc b/src/kernelDB.cc index af2483d..6b4d50d 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -291,6 +291,7 @@ bool kernelDB::addKernel(std::unique_ptr kernel) void kernelDB::ensureKernelLoaded(const std::string& name) { std::string canonical = getKernelName(name); + std::cerr << "[KDB] ensureKernelLoaded: canonical='" << canonical.substr(0, 80) << "'" << std::endl; std::string hsaco_path, elf_symbol; { std::shared_lock lock(mutex_); @@ -441,6 +442,7 @@ bool kernelDB::scanCodeObject(const std::string& co_file) bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName) { + std::cerr << "[KDB] scanCodeObjectForKernel: co=" << co_file << " kernel='" << kernelName.substr(0, 80) << "'" << std::endl; std::string strDisassembly; // Use targeted disassembly (--disassemble-symbols) to extract only the // requested kernel instead of disassembling the entire code object. @@ -454,15 +456,23 @@ bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::st { try { gotDisasm = getDisassemblyForSymbol(agent_, co_file, sym, strDisassembly); + } catch (const std::exception& e) { + std::cerr << "[KDB] getDisassemblyForSymbol threw: " << e.what() << std::endl; + gotDisasm = false; } catch (...) { + std::cerr << "[KDB] getDisassemblyForSymbol threw unknown exception" << std::endl; gotDisasm = false; } if (gotDisasm) + { + std::cerr << "[KDB] disassembly succeeded for symbol '" << sym.substr(0, 80) << "' (" << strDisassembly.size() << " bytes)" << std::endl; break; + } strDisassembly.clear(); } if (!gotDisasm) { + std::cerr << "[KDB] scanCodeObjectForKernel: no disassembly found" << std::endl; return false; } @@ -548,6 +558,8 @@ std::string kernelDB::extractKernelName(const std::string& line) parse_mode kernelDB::getLineType(std::string& line) { parse_mode result = BBLOCK; + if (line.empty()) + return result; auto it = line.begin(); if (*it == ':' || line.starts_with(".text:")) result = BEGIN; @@ -576,9 +588,13 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::map>::iterator it; do @@ -822,7 +838,13 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str skip = false; kernel = std::make_unique(demangledName); mit = markers.find(demangledName); - assert(mit != markers.end()); + if (mit == markers.end()) + { + skip = true; + current_kernel = nullptr; + current_block = nullptr; + break; + } current_kernel = kernel.get(); mode = BBLOCK; addKernel(std::move(kernel)); From af065a97573bb0c13d6a6bb91e9970c09180756f Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 05:17:22 -0700 Subject: [PATCH 07/12] Fix out-of-bounds access when parsing disassembly tokens The disassembly parser searched for '//' comment tokens by incrementing an index without bounds checking, causing a segfault when the disassembly line format didn't include address comments. Add bounds checks to all while-loops that search for '//' tokens in both getBlockMarkers and parseDisassemblyForKernel. Guard stoull calls with try/catch to handle unexpected formats. Co-Authored-By: Claude Opus 4.6 --- src/kernelDB.cc | 40 ++++++++++++++++++++++++++++++---------- 1 file changed, 30 insertions(+), 10 deletions(-) diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 6b4d50d..92159d4 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -625,15 +625,23 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::mapsecond.insert(base_addr + std::stoull(tmp[1], nullptr, 16)); } - else if (base_addr == 0) + else if (base_addr == 0 && tokens.size() > 1) { size_t i = 1; - while (tokens[i].find("//") == std::string::npos) + while (i < tokens.size() && tokens[i].find("//") == std::string::npos) i++; - std::string strAddress = tokens[++i]; - // remove the ending colon - strAddress.pop_back(); - base_addr = std::stoull(strAddress, nullptr, 16); + if (i + 1 < tokens.size()) + { + std::string strAddress = tokens[i + 1]; + // remove the ending colon + if (!strAddress.empty()) + strAddress.pop_back(); + try { + base_addr = std::stoull(strAddress, nullptr, 16); + } catch (...) { + base_addr = 0; + } + } } } }while(std::getline(in,line)); @@ -910,13 +918,25 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str inst.operands_.push_back(tokens[i]); } size_t i = 1; - while (tokens[i].find("//") == std::string::npos) + while (i < tokens.size() && tokens[i].find("//") == std::string::npos) { i++; } - std::string strAddress = tokens[++i]; - strAddress.pop_back(); - inst.address_ = std::stoull(strAddress, nullptr, 16); + if (i + 1 < tokens.size()) + { + std::string strAddress = tokens[i + 1]; + if (!strAddress.empty()) + strAddress.pop_back(); + try { + inst.address_ = std::stoull(strAddress, nullptr, 16); + } catch (...) { + inst.address_ = 0; + } + } + else + { + inst.address_ = 0; + } if (!blockCreated && mit != markers.end() && (mit->second.find(inst.address_) != mit->second.end())) { if (current_block) From 6164517f2050d4ebd8b1de61bfdbfbb353dfa191 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 1 Apr 2026 05:25:21 -0700 Subject: [PATCH 08/12] Fix all crash sites in disassembly parsing - getBlockMarkers: replace assert(tmp.size()==2) with defensive check for _cbranch_ token format; guard against uninitialized iterator when no KERNEL line has been seen yet - getLineType: guard *(--it) against single-character strings - parseDisassemblyForKernel: replace abort() with graceful skip when branch instruction seen without active kernel; guard null current_block before addInstruction call Co-Authored-By: Claude Opus 4.6 --- src/kernelDB.cc | 32 ++++++++++++++++++++------------ 1 file changed, 20 insertions(+), 12 deletions(-) diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 92159d4..e4c7d22 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -569,7 +569,7 @@ parse_mode kernelDB::getLineType(std::string& line) if (*it == ':') { // It's only a valid kernel if there are no spaces in lines that end with ':' - if(line.find_first_of(" ") == std::string::npos || *(--it) == '>') + if(line.find_first_of(" ") == std::string::npos || (line.size() > 1 && *(--it) == '>')) { if (line.ends_with("<.text>:")) result = BEGIN; @@ -596,7 +596,8 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::map>::iterator it; + std::map>::iterator it = markers.end(); + bool haveKernel = false; do { parse_mode mode = getLineType(line); @@ -611,8 +612,9 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::map tokens; split(line, tokens, " ", false); @@ -621,9 +623,15 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::map tmp; split(addr, tmp, "+", false); - assert(tmp.size() == 2); - tmp[1].pop_back(); - it->second.insert(base_addr + std::stoull(tmp[1], nullptr, 16)); + if (tmp.size() == 2 && !tmp[1].empty()) + { + tmp[1].pop_back(); + try { + it->second.insert(base_addr + std::stoull(tmp[1], nullptr, 16)); + } catch (...) { + // malformed address — skip this branch marker + } + } } else if (base_addr == 0 && tokens.size() > 1) { @@ -876,15 +884,14 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str trim(tokens[0]); if (isBranch(tokens[0])) { - if (current_kernel) + if (current_kernel && block) { current_kernel->addBlock(block_count, std::move(block)); } - else + else if (!current_kernel) { - std::cout << "Disassembly parsing error. Processing a branch instruction when there's not a kernel currently defined.\n"; - std::cout << line << std::endl; - abort(); + current_block = nullptr; + continue; } if (tokens[0].find("s_endpgm") != std::string::npos) { @@ -948,7 +955,8 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str current_block = block.get(); } inst.block_ = current_block; - current_block->addInstruction(inst); + if (current_block) + current_block->addInstruction(inst); if (inst.inst_ == "s_endpgm") { if (current_kernel && current_block) From 6d7ff8a957f0e21d92c2d01d6bd87e52f8c94b65 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 13 Apr 2026 07:31:56 -0700 Subject: [PATCH 09/12] Address reviewer must-fix items for lazy loading PR 1. Thread safety: ensureKernelLoaded() now uses a loading_kernels_ set and condition_variable so concurrent threads wait instead of both disassembling the same kernel. 2. Temp file leak: replace tmpnam() with mkstemp() in both getDisassembly() and getDisassemblyForSymbol(); unlink on all paths including failure. 3. Remove all unconditional [KDB] stderr debug prints. Co-Authored-By: Claude Opus 4.6 --- include/kernelDB.h | 3 +++ src/disassemble.cc | 35 +++++++++++++++++++---------------- src/kernelDB.cc | 41 ++++++++++++++++++++++------------------- 3 files changed, 44 insertions(+), 35 deletions(-) diff --git a/include/kernelDB.h b/include/kernelDB.h index cf61a43..b2cf036 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -254,6 +254,9 @@ class __attribute__((visibility("default"))) kernelDB { }; /// Lazy-loaded kernels: canonical name -> entry. Filled by addFile(..., lazy=true). std::map lazy_kernels_; + /// Kernels currently being loaded — prevents concurrent disassembly of the same kernel. + std::set loading_kernels_; + std::condition_variable_any loading_cv_; std::shared_mutex mutex_; }; diff --git a/src/disassemble.cc b/src/disassemble.cc index 7d65370..cf0b6d0 100644 --- a/src/disassemble.cc +++ b/src/disassemble.cc @@ -72,19 +72,19 @@ bool getDisassembly(hsa_agent_t agent, const std::string& fileName, std::string& ss << "--mcpu=" << name; parms.push_back(ss.str()); parms.push_back(fileName); - // Create a temporary file using tmpnam (note: tmpnam is not the most secure option) - char temp_filename[L_tmpnam]; - if (tmpnam(temp_filename) == nullptr) - throw std::runtime_error("Failed to generate temporary filename"); - else if (invokeProgram(disassembler, parms, temp_filename)) + char temp_template[] = "/tmp/kdb_dis_XXXXXX"; + int fd = mkstemp(temp_template); + if (fd < 0) + throw std::runtime_error("Failed to create temporary file"); + close(fd); + if (invokeProgram(disassembler, parms, temp_template)) { - // read file contents here - readFileToString(temp_filename, out); - unlink(temp_filename); + readFileToString(temp_template, out); + unlink(temp_template); return out.length() != 0; } - else - return false; + unlink(temp_template); + return false; } else return false; @@ -109,16 +109,19 @@ bool getDisassemblyForSymbol(hsa_agent_t agent, const std::string& fileName, parms.push_back("--disassemble-symbols=" + symbolName); parms.push_back(fileName); - char temp_filename[L_tmpnam]; - if (tmpnam(temp_filename) == nullptr) - throw std::runtime_error("Failed to generate temporary filename"); + char temp_template[] = "/tmp/kdb_dis_XXXXXX"; + int fd = mkstemp(temp_template); + if (fd < 0) + throw std::runtime_error("Failed to create temporary file"); + close(fd); - if (invokeProgram(disassembler, parms, temp_filename)) + if (invokeProgram(disassembler, parms, temp_template)) { - readFileToString(temp_filename, out); - unlink(temp_filename); + readFileToString(temp_template, out); + unlink(temp_template); return out.length() != 0; } + unlink(temp_template); return false; } diff --git a/src/kernelDB.cc b/src/kernelDB.cc index e4c7d22..d0e9132 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -291,26 +291,38 @@ bool kernelDB::addKernel(std::unique_ptr kernel) void kernelDB::ensureKernelLoaded(const std::string& name) { std::string canonical = getKernelName(name); - std::cerr << "[KDB] ensureKernelLoaded: canonical='" << canonical.substr(0, 80) << "'" << std::endl; std::string hsaco_path, elf_symbol; { - std::shared_lock lock(mutex_); + std::unique_lock lock(mutex_); + + // If another thread already loaded this kernel, nothing to do. auto it = lazy_kernels_.find(canonical); if (it == lazy_kernels_.end()) - { return; - } + + // If another thread is currently loading this kernel, wait for it. + while (loading_kernels_.count(canonical)) + loading_cv_.wait(lock); + + // Re-check: the other thread may have finished loading it. + it = lazy_kernels_.find(canonical); + if (it == lazy_kernels_.end()) + return; + hsaco_path = it->second.hsaco_path; elf_symbol = it->second.elf_symbol; + loading_kernels_.insert(canonical); } - if (!scanCodeObjectForKernel(hsaco_path, elf_symbol)) - { - return; - } + + bool ok = scanCodeObjectForKernel(hsaco_path, elf_symbol); + { std::unique_lock lock(mutex_); - lazy_kernels_.erase(canonical); + if (ok) + lazy_kernels_.erase(canonical); + loading_kernels_.erase(canonical); } + loading_cv_.notify_all(); } bool kernelDB::addFile(const std::string& name, hsa_agent_t agent, const std::string& strFilter, bool lazy) @@ -442,7 +454,6 @@ bool kernelDB::scanCodeObject(const std::string& co_file) bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName) { - std::cerr << "[KDB] scanCodeObjectForKernel: co=" << co_file << " kernel='" << kernelName.substr(0, 80) << "'" << std::endl; std::string strDisassembly; // Use targeted disassembly (--disassemble-symbols) to extract only the // requested kernel instead of disassembling the entire code object. @@ -456,25 +467,17 @@ bool kernelDB::scanCodeObjectForKernel(const std::string& co_file, const std::st { try { gotDisasm = getDisassemblyForSymbol(agent_, co_file, sym, strDisassembly); - } catch (const std::exception& e) { - std::cerr << "[KDB] getDisassemblyForSymbol threw: " << e.what() << std::endl; + } catch (const std::exception&) { gotDisasm = false; } catch (...) { - std::cerr << "[KDB] getDisassemblyForSymbol threw unknown exception" << std::endl; gotDisasm = false; } if (gotDisasm) - { - std::cerr << "[KDB] disassembly succeeded for symbol '" << sym.substr(0, 80) << "' (" << strDisassembly.size() << " bytes)" << std::endl; break; - } strDisassembly.clear(); } if (!gotDisasm) - { - std::cerr << "[KDB] scanCodeObjectForKernel: no disassembly found" << std::endl; return false; - } parseDisassemblyForKernel(strDisassembly, kernelName); From a8b87df7f81dda475c6c6189cd438b86ca1b7929 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 13 Apr 2026 07:34:43 -0700 Subject: [PATCH 10/12] Address reviewer should-fix items for lazy loading PR 4. Add test_lazy_loading.py: verifies lazy-loaded kernel discovery, has_kernel(), assembly, and per-line instructions match eager loading. 5. Move scanCodeObjectForKernel and parseDisassemblyForKernel to private. 6. Fail fast in ensureKernelLoaded(): always erase from lazy_kernels_ on failure so subsequent getKernel() calls get a clean error instead of retrying a doomed disassembly. 7. (Already done in must-fix #2: tmpnam -> mkstemp) 8. Eliminate ~170-line duplication: parseDisassembly() now delegates to parseDisassemblyForKernel(text, ""), using the safer bounds-checked version as the single implementation. Co-Authored-By: Claude Opus 4.6 --- include/kernelDB.h | 4 +- src/kernelDB.cc | 176 +++---------------------------------- tests/test_lazy_loading.py | 131 +++++++++++++++++++++++++++ 3 files changed, 146 insertions(+), 165 deletions(-) create mode 100644 tests/test_lazy_loading.py diff --git a/include/kernelDB.h b/include/kernelDB.h index b2cf036..9b871b7 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -226,10 +226,10 @@ class __attribute__((visibility("default"))) kernelDB { static void getElfSectionBits(const std::string &fileName, const std::string §ionName, size_t& offset, std::vector& sectionData ); std::vector getKernelArguments(const std::string& kernel_name, bool resolve_typedefs = false); bool scanCodeObject(const std::string& co_file); - bool scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName); - bool parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel); bool hasKernel(const std::string& name); private: + bool scanCodeObjectForKernel(const std::string& co_file, const std::string& kernelName); + bool parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel); /// Get kernel symbol names from a .hsaco ELF without disassembling (reads .symtab). static std::vector getKernelNamesFromElf(const std::string& fileName); /// If kernel is lazy-loaded, disassemble its code object and fill kernels_; then remove from lazy set. diff --git a/src/kernelDB.cc b/src/kernelDB.cc index d0e9132..2092f31 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -318,8 +318,10 @@ void kernelDB::ensureKernelLoaded(const std::string& name) { std::unique_lock lock(mutex_); - if (ok) - lazy_kernels_.erase(canonical); + // Always remove from lazy_kernels_: on success the kernel is now in + // kernels_; on failure we avoid retrying a doomed disassembly on every + // subsequent getKernel() call. + lazy_kernels_.erase(canonical); loading_kernels_.erase(canonical); } loading_cv_.notify_all(); @@ -661,159 +663,17 @@ void kernelDB::getBlockMarkers(const std::string& disassembly, std::map kernel; - CDNAKernel *current_kernel = nullptr; - std::unique_ptr block; - basicBlock *current_block = nullptr; - std::map> markers; - getBlockMarkers(text, markers); - //std::cout << "Found " << markers.size() << " in getBlockMarkers"; - std::map>::iterator mit; - bool bDoingKernels = false; - while(std::getline(in,line)) - { - bool blockCreated = false; - std::vector tokens; - mode = getLineType(line); - switch(mode) - { - case BEGIN: - mode = KERNEL; - block_count = 0; - break; - case KERNEL: - bDoingKernels = true; - strKernel = extractKernelName(line); - //strKernel = line.substr(0, line.length() - 1); - kernel = std::make_unique(demangleName(strKernel.c_str())); - mit = markers.find(demangleName(strKernel.c_str())); - assert(mit != markers.end()); - current_kernel = kernel.get(); - mode=BBLOCK; - addKernel(std::move(kernel)); - - break; - case BBLOCK: - if (!bDoingKernels) - break; - split(line, tokens, " ", false); - if (tokens.size()) - { - if (!current_block) - { - //std::cout << "Starting a new block:\n\t" << line << std::endl; - block = std::make_unique(); - block_count++; - current_block = block.get(); - blockCreated = true; - } - trim(tokens[0]); - if (isBranch(tokens[0])) - { - if (current_kernel) - current_kernel->addBlock(block_count, std::move(block)); - else - { - std::cout << "Disassembly parsing error. Processing a branch instruction when there's not a kernel currently defined.\n"; - std::cout << line << std::endl; - abort(); - } - if (tokens[0].find("s_endpgm") != std::string::npos) - { - blockCreated = true; - // std::cout << "New Block at endpgm\n\t" << line << std::endl; - block = std::make_unique(); - block_count++; - current_block = block.get(); - // Going to let this drop down to add the s_endpgrm instruction to the block. - // This is the only branch instruction we include - // This should always be the last block and last instruction in the kernel - } - else - { - current_block = nullptr; - continue; - } - } - std::vector inst_tokens; - instruction_t inst; - - // If there is more than one token and the first token contains underscores that means it's an instruction line - if (tokens.size() > 1 && tokens[0].find("_") != std::string::npos) - { - split(tokens[0], inst_tokens, "_", false); - if (inst_tokens.size() > 2 && (inst_tokens[1] == "load" || inst_tokens[1] == "store")) - { - inst.prefix_ = inst_tokens[0]; - inst.type_ = inst_tokens[1]; - inst.size_ = inst_tokens[2]; - } - inst.inst_ = tokens[0]; - inst.disassembly_ = line; - for (size_t i = 1; i < tokens.size() && tokens[i].find("//") == std::string::npos; i++) - inst.operands_.push_back(tokens[i]); - size_t i = 1; - while (tokens[i].find("//") == std::string::npos) - i++; - std::string strAddress = tokens[++i]; - // remove the ending colon - strAddress.pop_back(); - inst.address_ = std::stoull(strAddress, nullptr, 16); - if (!blockCreated && mit != markers.end() && (mit->second.find(inst.address_) != mit->second.end())) - { - // This is the first line of a new block - // So add the current block and create a new one - // Not all blocks end in branches, so when we come to - // an address identified as the target of a conditional branch - // we don't care if the current block ends with a branch instruction - // we just save it and create a new one. - //std::cout << "Starting block at " << strAddress << std::endl; - if (current_block) - current_kernel->addBlock(block_count, std::move(block)); - block = std::make_unique(); - block_count++; - current_block = block.get(); - } - inst.block_ = current_block; - current_block->addInstruction(inst); - if (inst.inst_ == "s_endpgm") - { - if (current_kernel && current_block) - { - current_kernel->addBlock(block_count, std::move(block)); - current_block = nullptr; - } - else - std::cerr << "Error parsing disassembly - s_endpgm without current kernel or block\n"; - } - } - } - break; - case BRANCH: - current_block = nullptr; - break; - default: - break; - } - } - - std::cout << "Marker Count: " << markers.size() << " Kernel Count: " << kernels_.size() << std::endl; - - return bReturn; + return parseDisassemblyForKernel(text, ""); } bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::string& targetKernel) { - // Normalize the target name so it can be compared against demangled canonical - // forms extracted from the disassembly output. The caller may pass a raw - // (mangled) ELF symbol name, so we demangle and canonicalize it here. - std::string targetCanonical = getKernelName(demangleName(targetKernel.c_str())); + // When targetKernel is non-empty, only parse that kernel and skip others. + // When empty, parse all kernels (full disassembly mode). + const bool filterByKernel = !targetKernel.empty(); + std::string targetCanonical; + if (filterByKernel) + targetCanonical = getKernelName(demangleName(targetKernel.c_str())); bool bReturn = true; std::istringstream in(text); @@ -829,7 +689,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str getBlockMarkers(text, markers); std::map>::iterator mit; bool bDoingKernels = false; - bool skip = false; // true when current kernel is not the target + bool skip = false; while(std::getline(in,line)) { bool blockCreated = false; @@ -847,7 +707,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str strKernel = extractKernelName(line); std::string demangledName = demangleName(strKernel.c_str()); std::string canonical = getKernelName(demangledName); - if (canonical != targetCanonical) + if (filterByKernel && canonical != targetCanonical) { skip = true; current_kernel = nullptr; @@ -871,9 +731,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str } case BBLOCK: if (!bDoingKernels || skip) - { break; - } split(line, tokens, " ", false); if (tokens.size()) { @@ -924,14 +782,10 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str inst.inst_ = tokens[0]; inst.disassembly_ = line; for (size_t i = 1; i < tokens.size() && tokens[i].find("//") == std::string::npos; i++) - { inst.operands_.push_back(tokens[i]); - } size_t i = 1; while (i < tokens.size() && tokens[i].find("//") == std::string::npos) - { i++; - } if (i + 1 < tokens.size()) { std::string strAddress = tokens[i + 1]; @@ -950,9 +804,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str if (!blockCreated && mit != markers.end() && (mit->second.find(inst.address_) != mit->second.end())) { if (current_block) - { current_kernel->addBlock(block_count, std::move(block)); - } block = std::make_unique(); block_count++; current_block = block.get(); @@ -968,9 +820,7 @@ bool kernelDB::parseDisassemblyForKernel(const std::string& text, const std::str current_block = nullptr; } else - { std::cerr << "Error parsing disassembly - s_endpgm without current kernel or block\n"; - } } } } diff --git a/tests/test_lazy_loading.py b/tests/test_lazy_loading.py new file mode 100644 index 0000000..4db9836 --- /dev/null +++ b/tests/test_lazy_loading.py @@ -0,0 +1,131 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + +""" +Tests for KernelDB lazy loading path: add_file(path, lazy=True) should +index kernel names without disassembly, then load on demand and produce +results identical to eager loading. +""" + +import subprocess +import tempfile +from pathlib import Path + +import pytest + +from conftest import requires_rocm + +kerneldb = pytest.importorskip("kerneldb", reason="kernelDB C++ extension not available") +KernelDB = kerneldb.KernelDB + +_HIP_SOURCE = r""" +#include + +__global__ void lazy_add(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] + b[idx]; + } +} + +__global__ void lazy_mul(float* a, float* b, float* c, int n) { + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < n) { + c[idx] = a[idx] * b[idx]; + } +} + +int main() { return 0; } +""" + +_binary_path = None + + +def _get_binary(): + global _binary_path + if _binary_path is not None: + return _binary_path + tmp = Path(tempfile.mkdtemp(prefix="kerneldb_lazy_test_")) + src = tmp / "lazy.cpp" + exe = tmp / "lazy" + src.write_text(_HIP_SOURCE) + r = subprocess.run(["hipcc", "-g", str(src), "-o", str(exe)], + capture_output=True, text=True) + if r.returncode != 0: + pytest.skip(f"hipcc compilation failed:\n{r.stderr}") + _binary_path = str(exe) + return _binary_path + + +def _find_kernel(kernels, fragment): + matches = [k for k in kernels if fragment in k] + if not matches: + pytest.skip(f"No kernel containing {fragment!r} found") + return matches[0] + + +@requires_rocm +def test_lazy_kernel_discovery(): + """Lazy add_file should discover the same kernels as eager loading.""" + path = _get_binary() + + eager = KernelDB(path) + eager_kernels = sorted(eager.get_kernels()) + + lazy = KernelDB(lazy=True) + lazy.add_file(path, lazy=True) + lazy_kernels = sorted(lazy.get_kernels()) + + assert lazy_kernels == eager_kernels + + +@requires_rocm +def test_lazy_has_kernel(): + """has_kernel() should return True for lazy-indexed kernels before loading.""" + path = _get_binary() + kdb = KernelDB(lazy=True) + kdb.add_file(path, lazy=True) + + kernels = kdb.get_kernels() + assert len(kernels) >= 2 + + for name in kernels: + assert kdb.has_kernel(name), f"has_kernel({name!r}) returned False" + + +@requires_rocm +def test_lazy_get_kernel_matches_eager(): + """Lazy-loaded kernel assembly should match eager-loaded assembly.""" + path = _get_binary() + + eager = KernelDB(path) + eager_kernel = eager.get_kernel(_find_kernel(eager.get_kernels(), "lazy_add")) + + lazy = KernelDB(lazy=True) + lazy.add_file(path, lazy=True) + lazy_kernel = lazy.get_kernel(_find_kernel(lazy.get_kernels(), "lazy_add")) + + assert lazy_kernel.assembly == eager_kernel.assembly + assert lazy_kernel.lines == eager_kernel.lines + + +@requires_rocm +def test_lazy_instructions_for_line(): + """Instructions retrieved via lazy loading should match eager loading.""" + path = _get_binary() + + eager = KernelDB(path) + eager_name = _find_kernel(eager.get_kernels(), "lazy_add") + eager_lines = eager.get_kernel_lines(eager_name) + + lazy = KernelDB(lazy=True) + lazy.add_file(path, lazy=True) + lazy_name = _find_kernel(lazy.get_kernels(), "lazy_add") + lazy_lines = lazy.get_kernel_lines(lazy_name) + + assert lazy_lines == eager_lines + + for line in eager_lines: + eager_insts = [i.disassembly for i in eager.get_instructions_for_line(eager_name, line)] + lazy_insts = [i.disassembly for i in lazy.get_instructions_for_line(lazy_name, line)] + assert lazy_insts == eager_insts, f"Mismatch at line {line}" From 967fff86db6e8d248cf85b3c27e30fc8cea2cc82 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 13 Apr 2026 07:35:19 -0700 Subject: [PATCH 11/12] Revert unrelated pyproject.toml license format change Split out per reviewer feedback (nice-to-have #10). Co-Authored-By: Claude Opus 4.6 --- pyproject.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyproject.toml b/pyproject.toml index f02b9d5..88f777b 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,7 +12,7 @@ description = "C++ library for querying data within CDNA kernel implementations, authors = [ { name = "Keith Lowery", email = "Keith.Lowery@amd.com" }, ] -license = { text = "MIT" } +license-files = ["LICENSE"] readme = "README.md" requires-python = ">=3.8" dependencies = [] From f87a0d9195884258b354a22810d9844774d03709 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 13 Apr 2026 07:48:40 -0700 Subject: [PATCH 12/12] Fix CI: replace condition_variable_any with condition_variable hipcc does not support std::condition_variable_any. Use a separate std::mutex + std::condition_variable for the loading sentinel instead of piggybacking on the shared_mutex. Add missing include. Co-Authored-By: Claude Opus 4.6 --- include/kernelDB.h | 4 +++- src/kernelDB.cc | 31 +++++++++++++++++-------------- 2 files changed, 20 insertions(+), 15 deletions(-) diff --git a/include/kernelDB.h b/include/kernelDB.h index 9b871b7..32360fd 100644 --- a/include/kernelDB.h +++ b/include/kernelDB.h @@ -53,6 +53,7 @@ THE SOFTWARE. #include #include #include +#include #include #include #include @@ -256,7 +257,8 @@ class __attribute__((visibility("default"))) kernelDB { std::map lazy_kernels_; /// Kernels currently being loaded — prevents concurrent disassembly of the same kernel. std::set loading_kernels_; - std::condition_variable_any loading_cv_; + std::mutex loading_mutex_; + std::condition_variable loading_cv_; std::shared_mutex mutex_; }; diff --git a/src/kernelDB.cc b/src/kernelDB.cc index 2092f31..f7d8b7f 100644 --- a/src/kernelDB.cc +++ b/src/kernelDB.cc @@ -292,36 +292,39 @@ void kernelDB::ensureKernelLoaded(const std::string& name) { std::string canonical = getKernelName(name); std::string hsaco_path, elf_symbol; - { - std::unique_lock lock(mutex_); - // If another thread already loaded this kernel, nothing to do. - auto it = lazy_kernels_.find(canonical); - if (it == lazy_kernels_.end()) - return; + // Use loading_mutex_ to serialize the check-and-claim of loading_kernels_. + { + std::unique_lock lk(loading_mutex_); // If another thread is currently loading this kernel, wait for it. while (loading_kernels_.count(canonical)) - loading_cv_.wait(lock); + loading_cv_.wait(lk); - // Re-check: the other thread may have finished loading it. - it = lazy_kernels_.find(canonical); - if (it == lazy_kernels_.end()) - return; + // Check lazy_kernels_ under shared_mutex. + { + std::shared_lock rlock(mutex_); + auto it = lazy_kernels_.find(canonical); + if (it == lazy_kernels_.end()) + return; // Already loaded (or was never lazy). + hsaco_path = it->second.hsaco_path; + elf_symbol = it->second.elf_symbol; + } - hsaco_path = it->second.hsaco_path; - elf_symbol = it->second.elf_symbol; loading_kernels_.insert(canonical); } bool ok = scanCodeObjectForKernel(hsaco_path, elf_symbol); { - std::unique_lock lock(mutex_); + std::unique_lock wlock(mutex_); // Always remove from lazy_kernels_: on success the kernel is now in // kernels_; on failure we avoid retrying a doomed disassembly on every // subsequent getKernel() call. lazy_kernels_.erase(canonical); + } + { + std::lock_guard lk(loading_mutex_); loading_kernels_.erase(canonical); } loading_cv_.notify_all();