Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs Cole Ramos
[email protected] AMD Austin, Texas, USA
arXiv:2605.03208v1 [cs.SE] 4 May 2026
Abstract Iterative GPU kernel tuning is bottlenecked by the scale of the applications that host the kernels. Rapid iteration requires isolating the kernel so it can be edited, recompiled, and validated without rebuilding the full application—but manual isolation requires reconstructing build flags, dispatch configuration, and runtime inputs by hand, so developers usually settle for slow in-place edits. We present Kerncap, an automated kernel extraction tool that intercepts dispatches at the HSA runtime for both HIP and Triton, bridging Triton’s JIT-only metadata into HSA-level capture via a lightweight Python compile-hook shim. Kerncap performs an address-space closure of all device memory—a virtual-addressfaithful snapshot that preserves embedded device pointers without DWARF metadata or pointer chasing—locates kernel sources, and emits self-contained reproducer projects. HIP reproducers use a Clang VFS overlay for source-level recompilation without modifying the original build system; Triton reproducers are tuning-pinned, binding the captured autotuner configuration into the artifact to preserve the JIT kernel’s numerical contract. Across six real-world HIP and Triton workloads spanning traditional HPC and ML domains on three AMD GPU architectures (CDNA2, CDNA3, RDNA3), Kerncap extracts and validates kernels from snapshots ranging from 152 MB to 30 GB—including a VAfaithful capture of vLLM’s Mixture-of-Experts weight pool reached through pointer indirection. On our llama.cpp case study, Kerncap’s edit-recompile-validate loop achieves a 13.6× speedup over the traditional workflow, reducing kernel isolation from a multihour process to a single command. The resulting reproducers also serve as a substrate for autotuning agents and LLM-driven kernel generators that need rapid, isolated evaluation of candidates.
1
Introduction
GPU kernel optimization is most effective when developers can iterate on a kernel in isolation: a tight edit-recompile-validate loop is far faster than rebuilding and rerunning an entire application. However, in practice, isolating a kernel is often the dominant cost in this workflow. As a concrete data point, on our llama.cpp case study a single full-application CMake rebuild takes 128 s—roughly 15× the application’s 8.3 s baseline runtime—and this cost is paid on every iteration. Consider a developer optimizing a single attention kernel within a large LLM inference pipeline containing hundreds of kernels. Although the target kernel may be small, extracting it requires navigating a multi-thousand-file codebase, resolving layers of templates and build system indirection, capturing gigabytes of deviceresident state, and reconstructing a runnable environment. A single mistake—missing a header dependency, misidentifying a dispatch
Keith Lowery
[email protected] AMD Austin, Texas, USA configuration, or capturing incomplete device memory—invalidates the reproducer. What should be a fast inner loop instead becomes a brittle, hours-long process. This difficulty arises because a kernel is not a self-contained artifact. Reproducing one requires simultaneously recovering three tightly coupled components, shown in Figure 1. Manual extraction must solve all three at once, which makes the process laborintensive and error-prone. 1. Definition. The compiled kernel binary and all transitive source dependencies— headers, translation units, and the build flags that produced them.
2. Runtime state. Grid and block dimensions, kernel arguments, and the contents of every device-memory region the kernel reads—including buffers reached indirectly through pointer arguments.
3. Environment. A build configuration and dispatch mechanism that faithfully replays the original execution with bit-identical semantics.
Figure 1: The three components of a GPU kernel reproducer. Each must be recovered in full for a replay to match the original dispatch, and the components are tightly coupled— runtime state references the definition, and the environment controls how both are reassembled. Kerncap automates this process end-to-end by capturing kernels at the moment of execution and reconstructing them into standalone, self-contained reproducer projects. With a single command: kerncap extract attn_fwd \ -- cmd " python bench . py " \ -- source - dir ./ src \ -- output ./ isolated / attn_fwd
Kerncap profiles the application, intercepts the target kernel dispatch, snapshots the complete device memory state, discovers all relevant source files, and generates a reproducer with built-in validation to ensure correctness. Relation to prior work. GPU kernel record-and-replay is not a new idea: NVIDIA’s Nsight Compute [9] provides kernel-, range-, and application-level replay modes for hardware counter collection, and CUPTI’s Checkpoint API (CUDA 11.5+) [8] saves and restores device state across replay passes. These tools, however, are closedsource, target NVIDIA hardware, and treat the captured state as an
Ramos, and Lowery
internal re-execution buffer rather than as an artifact the developer can read, edit, and recompile. Kerncap differs in three respects: it is open and AMD-native, it produces a self-contained reproducer that can be edited and rebuilt by the developer, and it captures the GPU address space verbatim rather than via explicit per-allocation save/restore hooks. Section 6 surveys the full competitive landscape. Contributions. Our open-source project1 currently targets AMD GPUs, supporting both HIP and Triton kernels. This paper makes the following contributions. • Unified HSA-level runtime interception. Automated kernel dispatch capture at the HSA level for both HIP and Triton applications, bridged for Triton by a lightweight Python compile-hook shim that supplies the just-in-time (JIT)-only metadata (user-visible kernel name, typed signature, constexprs) the HSA layer cannot see on its own. A single command-line interface (CLI) presents the unified capture path to the developer. • Address-space closure (VA-faithful capture). An address space is a closure over its pointer graph: by capturing every tracked GPU allocation at its original virtual address, the embedded pointer relationships (e.g., T** arguments) are preserved for free, without DWARF metadata or pointer chasing. We realize this with a full-address-space snapshot at the HSA level. • Automated source discovery. A multi-phase search algorithm that locates kernel definitions, traces #include dependencies, resolves translation units via compile_commands.json, and disambiguates template instantiation files using nm symbol lookup. • Self-contained reproducer generation. HIP reproducers use a Clang Virtual File System overlay that enables sourcelevel recompilation with the exact original compiler command. Triton reproducers are tuning-pinned: because JIT kernels carry an autotuner-dependent numerical contract that naive capture breaks, we bind the captured autotuner configuration into the reproducer to preserve it. • Validation framework. Smoke testing, byte-exact memory comparison for HIP variant validation, and tolerancebased comparison for Triton reproducers with NaN detection.
2 Background 2.1 GPU Kernel Execution on AMD AMD GPUs execute compute kernels through a layered software stack. At the lowest level, the Heterogeneous System Architecture (HSA) runtime manages kernel dispatch via AQL (Architected Queuing Language) packets submitted to hardware queues. Each kernel dispatch packet specifies the kernel code object handle, grid and workgroup dimensions, the kernel argument (kernarg) buffer address, and shared memory requirements. HSA tool interposition. AMD’s ROCProfiler-SDK [2] provides an LD_PRELOAD-based HSA interception mechanism—a registration API that exposes the live HsaApiTable, combined with 1 Kerncap is available at https://github.com/AMDResearch/intellikit
hsa_amd_queue_intercept_create for per-packet callbacks— that tools use to interpose on kernel dispatches and memory APIs. Kerncap builds on this mechanism; implementation details appear in Section 4.1. HIP. HIP (Heterogeneous-Compute Interface for Portability) is AMD’s C++ GPU programming model. Developers write __global__ functions compiled by hipcc (a wrapper around Clang) into HSACO (HSA Code Object) binaries. At runtime, HIP loads HSACOs into HSA executables, resolves kernel symbols, and dispatches them via the HSA runtime. Triton. OpenAI Triton is a Python-based GPU programming language [17]. Developers write kernel functions decorated with @triton.jit, which are JIT-compiled to native GPU code. The @triton.autotune decorator enables automatic selection among multiple tile-size configurations by benchmarking each and selecting the fastest. Importantly, different tile sizes change the floatingpoint accumulation order, causing significant numerical differences in half-precision computations. This means that a JIT-tuned kernel carries an implicit numerical contract between its tuning state and its outputs—a contract that any faithful reproducer must preserve. Empirically, on the FP16 attn_fwd kernel from flash-attn (B=2, H=16, S=4096, D=128) on MI300X, switching from the fastest autotune config to the second-fastest—only 7.7% slower—changes 11.3% of output elements with a maximum absolute error of 1.22 × 10−4 . The drift is attributed to the BLOCK_N tile, which reorders the softmax-denominator reduction across the K dimension. Kerncap preserves the contract by producing tuning-pinned reproducers (Section 4.4).
2.2
Manual Kernel Optimization Workflow
Since isolation is prohibitive, developers skip it entirely and operate in a coarser loop: (1) Profile: rank kernels by GPU execution time using rocprofv3 or similar tools. (2) Search: locate the kernel’s source code by tracing mangled names through headers, namespaces, and template instantiations—often across thousands of files. (3) Hypothesize: apply an optimization based on general heuristics (e.g., increase tile size, add prefetching, unroll an inner loop). (4) Rebuild: recompile the entire project or supporting library, which may take minutes. (5) Evaluate: re-run the full application and re-profile to see if the kernel improved. (6) Repeat: the optimization often fails to help—or makes things worse—because it was not tested against the kernel’s actual runtime inputs (problem sizes, memory layouts, argument values). The developer reverts and tries again. This loop is slow because every experiment touches the full build and the full application. There is no way to test a kernel change against the specific dispatch that was profiled without rebuilding and re-running everything. Writing a standalone harness for the kernel is possible—and is occasionally done for hot kernels in mature libraries—but it requires reconstructing the dispatch configuration,
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
input buffers, and build flags by hand for each new target. This upfront cost rarely pays off when the developer’s goal is exploratory optimization across dozens of kernels, so in practice they default to the outer rebuild-and-rerun loop.
3
Source Discovery. Locates the kernel definition and its translation unit: HIP via compile_commands.json and #include tracing, Triton via Python Abstract Syntax Tree (AST) walking for @triton.jit functions (Section 4.3). Reproducer Generation. Emits a self-contained project with a Clang VFS overlay (HIP) or a tuning-pinned Python replay script (Triton), preserving the original build-flag fidelity (Section 4.4). Validation. Replays the kernel and compares output: byteexact memory diff for HIP variants, tolerance-based numpy.allclose for Triton, or a smoke-test replay when no variant is supplied (Section 4.5).
Overview
The remainder of this paper presents Kerncap, which automates the workflow of Section 2.2 end-to-end. Kerncap is a commandline tool with a five-stage pipeline. At the user-facing layer, however, Kerncap exposes only three commands—kerncap profile, kerncap extract, and kerncap replay; the five stages describe the internal phases of an end-to-end run. Figure 2 illustrates the pipeline and the artifacts that flow between stages. kerncap profile 1. Profile rocprofv3 --kernel-trace
kerncap extract
target kernel
2. Capture runtime state HIP + Triton: LD_PRELOAD (rocprofiler-sdk) Triton: + Python compile-hook shim (metadata)
captured state 3. Source discovery HIP: compile_commands.json Triton: AST / Import tracing
source + deps 4. Reproducer generation HIP: Clang VFS Triton: Replay script
Unified HSA-level capture, language-specific reproducers. HIP and Triton kernels operate at fundamentally different abstraction levels: HIP kernels are compiled ahead-of-time to binary HSACOs and dispatched via HSA, while Triton kernels exist as Python objects JITcompiled at runtime. Kerncap converges both onto the same HSAlevel capture pipeline—bridging Triton’s JIT-only metadata with a lightweight Python compile-hook shim—and then diverges back into language-specific reproducer generation (VFS overlay for HIP, templated Python script for Triton). The result is a single kerncap extract CLI command that auto-detects the kernel language. Detection is a lightweight scan of the user-supplied --source-dir: if any .py file contains a @triton.jit or @triton.autotune decorator on a function whose name matches the target kernel, the Triton path is selected; otherwise the kernel is treated as HIP. An explicit --language user flag can override this heuristic when needed.
4
Design and Implementation
This section describes the five pipeline stages in detail. The implementation is a hybrid C++/Python system: the HSA interception library and replay binary are written in C++ for direct access to the HSA API table and low-level memory operations, while the CLI, source discovery, reproducer generation, Triton capture, and validation are implemented in Python.
reproducer
kerncap replay
4.1 5. Replay & validation HIP: Byte comp. Triton: allclose
Figure 2: The Kerncap workflow. The tool abstracts low-level instrumentation into three high-level commands, supporting both HIP and Triton backends. Profile. Ranks kernels by GPU time via rocprofv3 --kernel-trace --stats; included for workflow completeness, often skipped when the target kernel is already known. Capture. A shared library (libkerncap.so) loaded via LD_PRELOAD intercepts HSA dispatches for both HIP and Triton; a lightweight Python compile-hook shim supplies the JIT-only metadata (user-visible names, typed signatures, constexprs, autotuner configurations) that Triton requires, indexed by HSACO SHA-256 in name_map.json (Section 4.1).
Runtime Interception
4.1.1 HIP Path: HSA API Table Hooking. The HIP capture library, libkerncap.so, is loaded via LD_PRELOAD and registers with the rocprofiler-sdk framework. It exports a rocprofiler_configure entry point that calls rocprofiler_at_intercept_table_registration to install a callback. When the HSA runtime initializes, this callback receives the live HsaApiTable—a struct containing function pointers for the entire HSA API surface. Kerncap saves a copy of the original function pointers and replaces entries in the live table with its own implementations: (1) Queue interception. hsa_queue_create is replaced to call hsa_amd_queue_intercept_create, which creates an intercept queue, followed by hsa_amd_queue_intercept_register to install a per-packet callback (on_submit_packet). (2) Memory tracking. hsa_amd_memory_pool_allocate, hsa_memory_allocate, and the virtual-memory
Ramos, and Lowery
(VMEM) APIs (hsa_amd_vmem_address_reserve, hsa_amd_vmem_map, and their free/unmap counterparts) are hooked to maintain a map from device pointer to allocation size. (3) Symbol tracking. hsa_executable_get_symbol_by_name and hsa_executable_symbol_get_info are hooked to associate kernel object handles with their mangled symbol names. (4) Code object capture. hsa_code_object_reader_create_from_memory and hsa_executable_load_agent_code_object are hooked to intercept HSACO binary blobs as they are loaded into the runtime. The association from kernel object handle to HSACO blob is built lazily: when the runtime queries HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, the tool follows the chain kernel_object → symbol → executable → blob. Figure 3 illustrates the hook architecture. signal interposition wait → snapshot → release
Kerncap HSA calls ptr → size memory HIP Application obj → name symbols
forward
GPU
obj → blob HSACO capture
dispatch.json + *.bin
Figure 3: HSA hook architecture. Kerncap interposes on the HSA API table, maintaining three data structures: a pointerto-size map for memory tracking, a symbol-to-name map for kernel identification, and a blob store for HSACO code objects. All calls are forwarded to the original HSA implementation. Signal interposition for post-execution capture. When the packet callback identifies a target kernel dispatch, it cannot simply capture state immediately—the kernel has not yet executed, and output buffers contain stale data. Kerncap replaces the dispatch packet’s completion signal with a fresh signal, forwards the modified packet to the hardware, and blocks on the new signal. Once the kernel completes, Kerncap performs the memory snapshot, then decrements the original signal so the application proceeds normally. 4.1.2 Triton Path: HSA-Level Capture with Compile-Hook Shim. Triton kernels ultimately dispatch through the HSA runtime, so the same libkerncap.so mechanism that intercepts HIP dispatches captures Triton dispatches as well. What HSA does not see is the Triton-level metadata the reproducer needs: the user-visible kernel name, the typed Python signature, the values of tl.constexpr arguments, and the tensor dtypes and strides that exist only as PyTorch metadata. We bridge this gap with a lightweight Python compile-hook shim that runs alongside the HSA interception path: the shim records Triton-level metadata at compile time into name_map.json (indexed by HSACO SHA-256), and the HSA hook
cross-references that file at dispatch time to reassemble the Tritonlevel call from the binary kernarg buffer. Compile-hook shim. A sitecustomize.py module installs a hook on Triton’s JIT compilation path at interpreter startup. Because Python’s site module imports sitecustomize in every interpreter, the hook propagates to child processes spawned via multiprocessing.spawn—including vLLM’s EngineCore workers. When Triton compiles a kernel, the shim records the SHA-256 of the emitted HSACO together with the kernel’s user-visible name, typed signature, constexpr bindings, and per-argument tensor metadata into a shared name_map.json. HSA-level kernarg recovery. At dispatch time the HSA hook sees a kernel object handle and a binary kernarg buffer, but no typed layout. A C++ kernarg-metadata parser, invoked at executable-load time via amd_comgr, reads the AMDGPU code-object metadata blob and extracts each slot’s offset, size, value_kind, and type. Crossreferencing this typed layout against name_map.json (indexed by HSACO SHA-256) reassembles the original Triton call: mangled symbols resolve back to user-visible names, binary kernarg bytes decode to typed pointer and scalar arguments, and constexprs recover with their compile-time values. The resulting dispatch.json and memory_regions.json are format-identical to HIP captures, so the VA-faithful memory snapshot mechanism (Section 4.2) applies uniformly. Autotuner interception. When a kernel uses @triton.autotune, the compile-hook shim also records which keyword arguments originate from autotuner configurations (e.g., BLOCK_M, num_warps) and persists the winning configuration to name_map.json, pinning tile sizes, num_warps, and num_stages for deterministic replay. Early host termination. For long-running host applications, completing the full run after the target kernel has already been captured wastes wall time and risks resource contention with subsequent kerncap invocations. Both capture paths therefore drop a capture_complete sentinel file as the final act of artifact serialization. The kerncap extract CLI runs a watchdog thread that polls for this sentinel and SIGTERM/SIGKILLs the child process group the moment it appears, regardless of whether the host would otherwise have terminated in seconds or hours. This is the mechanism behind the Triton overhead footnote in Table 3.
4.2
Memory Capture
4.2.1 HIP Path: VA-Faithful Device Memory Snapshot. After the target kernel completes execution, Kerncap snapshots all tracked device memory regions—not just the kernel’s direct arguments. The snapshot_all_tracked_memory function iterates over the pointer-to-size map and streams each region to memory/region_{base_addr}.bin via hsa_memory_copy in fixed-size chunks (default 64 MiB, tunable via KERNCAP_SNAPSHOT_CHUNK_BYTES). Chunked streaming bounds the per-region host-memory footprint, which matters for multi-GiB allocations such as vLLM’s KV-cache slabs where a per-region buffer would otherwise spike host memory consumption.
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
A companion file, memory_regions.json, records each region’s base address, size, allocation type (pool vs. VMEM), and whether it contains the kernarg buffer. The kernarg segment itself is captured separately (kernarg.bin) with its exact size queried via HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE. Handling embedded device pointers via address-space closure. Many GPU kernels use indirect pointers: a kernel argument of type T** points to a device buffer that itself contains pointers to other device buffers. Traditional approaches would require parsing DWARF debug information or recursively chasing pointers through device memory. Kerncap sidesteps this entirely with what we call an address-space closure: an address space is a closure over its pointer graph, so capturing every tracked allocation at its original virtual address captures the graph for free. The replay binary restores memory at the same virtual addresses, so embedded pointers remain valid without any interpretation, regardless of how deeply nested they are. Bounds of the closure. Address-space closure is a device-side guarantee. Its assumptions break in three identifiable cases. (1) Hostresident pointers in kernel arguments—unified-memory pointers or pointers into mmap’d host files—are not captured by the devicememory snapshot and cannot be restored by device VMEM allocation. (2) Address-layout drift across processes: the closure guarantees device VAs are reproducible because the replay can request exact addresses via hsa_amd_vmem_address_reserve, but host-side pointers affected by address space layout randomization (ASLR) and passed in kernargs would not survive a fresh process. (3) Kerneltime pointer arithmetic that escapes the captured set—e.g., a kernel that dereferences a device pointer the runtime never tracked—is by construction outside the closure. Section 7 discusses the broader implications for host-side state. Module-variable capture. The pointer-to-size map tracks only allocations routed through HSA’s runtime memory APIs. It does not cover __constant__ memory populated at executable-load time via hipMemcpyToSymbol, which appears as HSA_SYMBOL_KIND_VARIABLE symbols embedded in the loaded executable. Portability layers—notably Kokkos’s kokkos_impl_hip_constant_memory_buffer—rely on this path for runtime constants that kernels then dereference, so ignoring it causes replays to fault on NULL View pointers. Kerncap enumerates these variable symbols during capture, reads each one’s contents into a per-capture module_vars/ directory, and restores them in replay after VMEM allocation but before kernel dispatch. When a process loads many executables that export the same symbol name—each Kokkos kernel carries its own kokkos_impl_hip_constant_memory_buffer— Kerncap disambiguates the restore by matching the executable’s SHA-256 against the HSACO recorded in the capture, ensuring the correct blob is restored even when hsa_executable_get_symbol_by_name would otherwise alias them. Ordering and crash safety. Kerncap writes memory_regions.json and dispatch.json before performing the device-to-host copies. If an application frees a device buffer between kernel completion and the snapshot (a race condition inherent to interception tools), the hsa_memory_copy for that region will fail, but the metadata
files and all successfully-copied regions remain intact. The HSACO binary is also saved before the memory snapshot for the same reason. 4.2.2 Triton Path: Unified with HIP. The HSA-based Triton backend (Section 4.1.2) produces memory_regions.json and region_{base}.bin files in the same format as HIP captures, and therefore inherits the same VA-faithful snapshot, address-space closure, chunked streaming, and module-variable capture mechanisms described above.
4.3
Source Discovery
4.3.1 HIP Path: Compile Database and Include Tracking. The HIP source finder takes three inputs—the kernel’s demangled name, its mangled symbol from dispatch.json, and a user source directory— and resolves the editable source plus its compile-unit translation unit. Figure 4 shows the fallback order. Kernel name + mangled symbol + source dir
yes
compile_commands.json + debug info available?
DWARF path nm locates obj; llvm-dwarfdump reads line table
no
grep fallback 2-pass: strict __global__ → loose Translation-unit resolution (nm disambiguation for template-instantiated TUs)
Editable source + TU + compile flags
Figure 4: HIP source-discovery decision tree. DWARF-based discovery is the primary path when debug info and a compile database are both available; grep-based search is the fallback. nm disambiguation resolves multi-candidate translation units (e.g., template instantiation files) in both paths. Inputs and base-name extraction. A demangled C++ kernel name often carries namespaces, template parameters, and (for truncated profiler output) unbalanced template brackets. The _extract_base_name routine strips these to obtain the unqualified function name used for grep-based matching, handling cases such as ck::GridwiseGemm<. . . >::Run → Run and the common llama.cpp case mul_mat_vec_q<(ggml_type)39. DWARF-first, grep-fallback. When a mangled symbol and a compile_commands.json are both available, Kerncap runs nm across the listed object files to locate the one containing the mangled kernel symbol, then reads that object’s DWARF line tables (via llvm-dwarfdump, with a readelf fallback) to extract all source files that contributed to the translation unit, filtered by the user’s --source-dir. This path is the only one that works for frameworkgenerated kernels (e.g., Kokkos) whose __global__ qualifier lives in framework headers rather than user code [13]. If debug info is
Ramos, and Lowery
absent or the compile database is missing, Kerncap falls back to a two-pass grep search under the source directory: a strict pass matching the __global__ function definition (to avoid driver files that merely reference the kernel), and a loose pass on the base name if the strict pass returns nothing. #include "..." directives are then recursively traced from the discovered file up to five levels deep to collect header dependencies. Translation-unit resolution. Kernels defined in headers compile through a separate .cu/.cpp translation unit. Both discovery paths resolve this the same way: scan compile_commands.json entries whose source #includes the kernel header (or, if the database is absent, grep .cu files for the include). When multiple candidates match—routine for template-instantiation file sets like llama.cpp’s mmq-instance-*.cu—nm disambiguates by selecting the candidate whose compiled object contains the exact mangled symbol from dispatch.json. Compile flags (-D, -I) come from the matched entry; when no entry is found, Kerncap infers defines by scanning source files for HIP/ROCm/AMD-guarded #ifdef blocks. Empirical incidence. On the two HIP workloads in our evaluation (llama.cpp and LAMMPS), the DWARF path succeeded for both, and nm disambiguation fired for both (llama.cpp’s templateinstantiated mul_mat_vec_q variants and LAMMPS’s Kokkosexpanded TagPairEAMKernelC). The grep fallback was not loadbearing in our measurements, but is retained for projects that ship without debug info or a compile database. 4.3.2 Triton Path: AST and Import Tracing. The Triton source finder operates on Python ASTs: (1) Parse every .py file in the source directory. (2) Walk the AST for FunctionDef nodes decorated with @triton.jit or @triton.autotune. (3) Match the kernel name against the function name (supporting substring matching for flexibility). (4) Trace ImportFrom nodes to find helper modules, supporting both relative imports (from .common import ...) and absolute imports. (5) Detect package structure (__init__.py) to preserve import hierarchies.
4.4
Reproducer Generation
4.4.1 HIP Reproducers. The HIP reproducer generator creates a self-contained project directory with the structure shown in Figure 5.
isolated/my_kernel/ capture/ dispatch.json kernarg.bin kernel.hsaco memory_regions.json memory/ region_7f8a00000000.bin region_7f8a04000000.bin ... kernel_variant.cpp deps/ common.hpp utils.cuh vfs.yaml Makefile
Figure 5: Structure of a generated HIP reproducer project. The key components are: Capture data. The capture/ directory contains the complete VA-faithful snapshot: dispatch metadata, kernarg buffer, HSACO binary, and all device memory regions. The replay binary (kerncap-replay) reads this directory directly. Editable source. The main translation unit is copied as kernel_variant.cpp, and all traced header dependencies are flattened into deps/ with collision handling (directory-prefixed names when two headers share a basename). VFS overlay. The Clang Virtual File System overlay (vfs.yaml) is the mechanism that enables source-level recompilation without modifying the original build system. It maps each local file copy back to its original filesystem path: {" version ": 0, " roots ": [ {" type ": " directory " , " name ": "/ original / src / kernels " , " contents ": [ {" type ": " file " , " name ": " gemm . cuh " , " external - contents ": "/ isolated / deps / gemm . cuh "} ]} ]}
When the Makefile’s recompile target runs, it invokes the exact original compiler command from compile_commands.json with three additional flags: -ivfsoverlay vfs.yaml (to substitute edited files), --cuda-device-only and --no-gpu-bundle-output (to produce a raw HSACO instead of a fat binary). This ensures 100% flag and dependency fidelity: the recompiled HSACO uses the same optimization level, architecture target, and include paths as the original build. Makefile. The generated Makefile provides four key targets: run (replay the captured kernel), recompile (rebuild the HSACO from edited source), run-variant (replay with the recompiled HSACO), and validate-variant (compare variant output against baseline).
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
4.4.2 Triton Reproducers. Triton reproducers are generated from a Jinja2 template that produces a standalone reproducer.py script. The script: (1) Loads captured tensor data from binary files using numpy.fromfile, converting to the original torch dtype and shape. (2) Imports the kernel function from the copied source module. (3) Launches the kernel with the captured grid dimensions and arguments. (4) Saves output tensors for validation. Tuning-pinned reproducers. A JIT-compiled Triton kernel carries an implicit numerical contract between its tuning state (tile sizes, warp/stage counts) and its outputs (Section 2.1). Naive replay breaks this contract: re-running an extracted kernel under the autotuner can silently select a different configuration and produce numerically different outputs. We preserve the contract by producing tuning-pinned reproducers. If the captured kernel used @triton.autotune, the reproducer bypasses the autotuner entirely; instead of calling the autotuner wrapper, it invokes kernel.fn[grid](**config) directly with the captured winning configuration (tile sizes, num_warps, num_stages), binding the tuning state into the reproducer artifact. Package-aware source copying. If the kernel module lives inside a Python package (directory contains __init__.py), the generator copies the entire package directory rather than individual files, preserving relative import chains.
4.5
Replay and Validation
4.5.1 VA-Faithful HSA Replay. The replay binary (kerncap-replay) restores the captured device memory state and re-dispatches the kernel using raw HSA APIs, entirely bypassing the HIP runtime. The replay proceeds in six stages: (1) Parse metadata. Read memory_regions.json to determine the set of virtual address ranges that must be restored, and dispatch.json for grid dimensions, kernarg size, and the mangled kernel symbol name. Before hsa_init(), issue (2) Pre-mmap. mmap(MAP_FIXED_NOREPLACE) calls for each captured region. This reserves the virtual address ranges so the HSA runtime’s SVM (Shared Virtual Memory) aperture initialization does not claim them. (3) HSA initialization. Call hsa_init(), then munmap the pre-reserved ranges. (4) VMEM allocation. For each region, call hsa_amd_vmem_address_reserve with the exact captured base address, create a backing handle via hsa_amd_vmem_handle_create, and map it with hsa_amd_vmem_map. If the reserve returns a different address than requested, the replay aborts—VA faithfulness is a hard requirement. (5) Memory restore and dispatch. Copy captured data into each restored region. Load the HSACO (or an override via --hsaco), resolve the kernel symbol by mangled name, allocate a kernarg buffer, fill it from kernarg.bin, and submit an AQL dispatch packet.
(6) Output dump. If --dump-output is specified, copy all regions back to host after kernel completion and write them to output/region_{base}.bin for comparison. The replay supports multi-iteration benchmarking with optional memory re-copy between iterations (--no-recopy for stateful measurement). Robustness. We have validated the replay pipeline on ROCm ≥ 7.0, where the rocprofiler-sdk interception framework is stable. The Stage 3 relocation-abort path was not observed to trigger across our evaluation. Should a future driver invalidate the pre-reservation assumption, the replay aborts loudly with a descriptive error rather than silently corrupting the captured pointer graph. 4.5.2 Replay Validation. The validator implements three comparison strategies: Smoke test (HIP baseline). When no variant HSACO is provided, validation simply confirms that the captured kernel replays without crashing. This serves as a basic sanity check that the capture data is complete and the address space was restored correctly. Byte-exact comparison (HIP variant). When --hsaco provides a recompiled HSACO, the validator runs two full replays: one with the captured HSACO (baseline) and one with the variant HSACO. Both replays use --dump-output to snapshot post-execution memory. The validator then compares every output region byte-for-byte, reporting the number of differing bytes and their percentage of total region size. Tolerance-based comparison (Triton). For Triton reproducers, the validator runs reproducer.py and compares the output tensors against captured reference data using numpy.allclose with configurable absolute (atol) and relative (rtol) tolerances. The validator detects and reports NaN values explicitly, noting that they typically indicate uninitialized device memory, half-precision overflow, or buffer size misinterpretation. Closing the loop. Returning to Figure 1: the definition is recovered by HSACO capture (Section 4.1) together with DWARF-based source discovery (Section 4.3) and VFS-flattened reproducer assembly (Section 4.4); the runtime state by address-space closure— the VA-faithful device-memory snapshot that preserves embedded pointer graphs without DWARF or pointer chasing (Section 4.2); and the environment by exact-flag single-file recompile via Clang VFS (Section 4.4) and VA-faithful HSA replay (Section 4.5). The tight coupling between components—particularly between runtime-state capture and environment replay—is why Kerncap treats capture and replay as a single mechanism rather than two independent tools.
5
Results
We evaluate Kerncap along three research questions: • RQ1 (Correctness and breadth). Does Kerncap successfully extract and validate kernels across a diverse set of real-world workloads, spanning both HIP and Triton and varying application complexity? • RQ2 (Overhead). What is the wall-clock cost imposed on the host application during capture, and how does it
Ramos, and Lowery
decompose into per-dispatch interception cost versus onetime memory snapshot cost? • RQ3 (Iteration speedup). How much does the isolated edit-recompile-validate loop accelerate kernel optimization compared to the traditional full-application workflow?
5.1
Experimental Setup
We evaluate Kerncap across three AMD GPU architectures spanning AMD’s data-center (CDNA) and consumer (RDNA) GPU families: AMD Instinct™ MI300X (gfx942, CDNA3), AMD Instinct™ MI210 (gfx90a, CDNA2), and AMD Radeon™ PRO W7900 (gfx1100, RDNA3). The CDNA hosts use AMD EPYC™ 9684X CPUs; the RDNA host uses an AMD Ryzen™ Threadripper PRO 5975WX. All systems run ROCm 7.2.0 on RHEL 9.6. Table 1 summarizes the six benchmark workloads. Table 1: Benchmark workloads used for evaluation, spanning three orthogonal axes: compiled vs. JIT (HIP, Triton, Inductor), authored vs. framework-generated (hand-written, Kokkos, Tensile, Inductor), and snapshot sizes from 152 MB to 30 GB. Workload
Lang.
Target kernel
Driver / Input
llama.cpp
HIP
mul_mat_vec_q
LAMMPS
HIP
TagPairEAMKernelC
llama-bench, GPT-OSS 20B MXFP4 lmp (Kokkos), EAM, 83 grid
rocBLAS GEMM
HIP
Cijk_Ailk_Bljk_HHS_BH
Flash Att. 2
Triton
attn_fwd
vLLM
Triton
fused_moe_kernel
torch.compile
Triton
triton_poi_fused_relu_0
rocblas-bench, sgemm 40963 sample_attn_layer.py, FP16 vllm bench latency, Qwen1.5-MoE-A2.7B torch.compile micro-bench, fp16 matmul
The six workloads span three orthogonal axes that the design space requires us to defend: compiled vs. JIT (HIP ahead-of-time vs. Triton runtime vs. torch.compile’s Inductor backend on-demand); authored vs. framework-generated (hand-written CUDA-style kernels vs. Kokkos template instantiation vs. Tensile/Inductor codegen); and small vs. large memory footprint (152 MB to 30 GB snapshots). Two workloads (rocBLAS GEMM, torch.compile) deliberately have no recoverable source, exercising the HSACO-only fallback path described in Section 4.3.
5.2
Extraction and Validation
For each workload we run the full Kerncap pipeline (kerncap extract followed by kerncap validate) on each available architecture and record the number of device memory regions captured, the total snapshot size, and the end-to-end extraction wall-clock time. Table 2 presents the results across the three GPUs. Sourcefile discovery is workload-stable (6 files for llama.cpp, 10–11 for LAMMPS, 2 for the hand-authored Triton workloads, 0 for the nosource workloads) and is omitted from the table for compactness.
Table 2: Cross-architecture extraction results. Workload
Arch
Reg.
Snap. (MB)
Time (s)
llama.cpp
gfx942 gfx90a gfx1100
21 21 20
12,233 12,126 11,520
15.6 20.9 18.8
LAMMPS
gfx942 gfx90a gfx1100
84 84 83
8,449 8,421 8,415
25.0 28.2 47.7
rocBLAS GEMM
gfx942 gfx90a gfx1100
10 10 10
248 152 152
9.0 11.2 4.4
Flash Att. 2
gfx942 gfx90a
13 13
178 178
9.1 13.2
vLLM
gfx942 gfx90a
185 186
30,074 29,301
60.9 85.2
torch.compile
gfx942 gfx90a
12 12
227 227
9.4 13.5
Reg. = device memory regions captured; Snap. = total snapshot size; Time = end-to-end wall-clock cost of kerncap extract, including post-capture reproducer generation. The instrumented host-app run alone is reported in Table 3; the difference (largest for LAMMPS at ∼10 s) is dominated by source-discovery and reproducer-assembly cost. Triton workloads tested on gfx942 and gfx90a only.
Kerncap extracts and validates kernels from every workload on every architecture for which it was tested. Run-to-run drift in the captured-region count (e.g., LAMMPS at 84 vs. 83 regions across archs) reflects natural variation in the Kokkos memory pool’s slab allocations between runs, not a capture-path discrepancy. Snapshot sizes are otherwise stable across architectures within a few percent, confirming that the HSA-level capture path enumerates the same VAs regardless of underlying ISA family. HIP validation. The llama.cpp reproducer passes full validation on all three architectures: the captured kernel replays without error, and recompiling the reproducer’s source without modification (make recompile) produces byte-identical output across all 21 compared memory regions on every arch. This confirms endto-end source-level fidelity through the VFS overlay mechanism on both CDNA and RDNA targets [5]. On gfx1100, the kernel replays correctly with the wave32 codegen path (Block: 32×1×1, vs. wave64’s 32×2 on CDNA), demonstrating that the capture metadata and replay harness are wavefront-width agnostic. LAMMPS similarly passes full validation across all three architectures: the isolated TagPairEAMKernelC kernel replays byte-identically with 83–84 restored memory regions spanning ∼8.4 GB of device state per arch [16]. HSACO-only path (no source). Two workloads—rocBLAS GEMM (Tensile codegen) and torch.compile (Inductor JIT)—deliberately exercise the HSACO-only reproducer path, in which source recovery is out of scope because the kernel’s source representation is either an internal codegen template (Tensile) or a runtime-emitted JIT artifact that does not persist in a stable on-disk location (Inductor). In both cases, Kerncap gracefully degrades: the captured HSACO replays correctly under kerncap replay, and kerncap validate confirms bit-exact replay against the captured baseline. The make recompile edit-loop is unavailable for these workloads by design,
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
but the reproducer remains useful for hardware-counter profiling, autotuner experimentation, and input-perturbation studies— capabilities no other tool currently provides for Tensile or Inductoremitted kernels. On gfx1100, rocBLAS GEMM extracts and validates in 4.4 s, confirming that the HSACO-only path is portable across ISA families. Triton validation. Both hand-authored Triton workloads (Flash Attention and vLLM) pass element-wise validation with zero error (max_error=0.0) on the architectures tested. The reproducers pin the autotuner configuration captured at runtime, bypassing re-tuning and ensuring deterministic replay [4, 6]. vLLM’s fused_moe_kernel is a noteworthy demonstration of VA-faithful capture: the kernel dereferences indirection layers (the per-expert weight tables in the MoE routing path), and Kerncap’s HSA-based backend captures the full device-memory closure those pointers reach—185 regions totaling ∼30 GB, dominated by 24×702 MB gate/up-projection tensors and 24×330 MB down-projection tensors (∼24 GB of expert weights, ∼5 GB activations/workspace)—enabling bit-exact replay of kernels that argument-only capture cannot reproduce.
5.3
Capture Overhead
A naïve overhead measurement—comparing total wall-clock time with and without instrumentation—conflates two fundamentally different costs: (1) the interception tax, a continuous per-dispatch callback cost that scales with kernel count, and (2) the capture cost, a one-time device-memory snapshot that scales with memory footprint, not application duration. Reporting a single overhead ratio lets workload duration confound the result: for comparable capture costs (∼7–14 s on llama.cpp across archs), a 10-second baseline yields a 1.9–2.6× wall-clock ratio, while a multi-minute workload such as vLLM sees the same fixed cost amortized to roughly 11% overhead. We therefore decompose Kerncap’s overhead into these two components. For each configuration we perform 𝑁 =10 measured runs after 1 warmup run (discarded) and report median ± std. dev. For HIP workloads, we measure an interception-only configuration by loading libkerncap.so via LD_PRELOAD while targeting a nonexistent kernel: all HSA hooks and per-dispatch callbacks execute, but no capture occurs. The full capture configuration targets the real kernel, triggering the one-time memory snapshot. For Triton workloads, the Python-level JITFunction.run monkey-patch adds negligible per-call overhead (a string comparison) without capture, so interception-only is omitted. Table 3 presents the wall-clock decomposition.
Table 3: Wall-clock overhead decomposition on gfx942 (𝑁 =10, 1 warmup). Intercept. is the ratio of interception-only wall time to Base. HIP Full cap. runs the host application to completion whereas Triton terminates once the target dispatch is captured. Full cap. measures the instrumented host-app run only; end-to-end kerncap extract time including reproducer assembly is reported in Table 2. Workload
Base (s)
Intercept. Full cap. Snap. Cap. ratio (s) (MB) cost (s)
HIP (LD_PRELOAD instrumentation) llama.cpp 8.30 ± 0.04 1.06× LAMMPS 9.96 ± 0.02 1.04× rocBLAS GEMM 3.90 ± 0.03 2.04×
15.8 15.3 8.2
12,233 8,449 248
7.0 4.9 0.2
9.1 60.8‡ 9.4
178 30,074 227
∼4.8† ∼18‡ 0.4†
Triton (Python-level instrumentation) Flash Att. 2 vLLM torch.compile
4.29 ± 0.03 163.8 ± 4.9 9.01 ± 0.09
— — —
† Cap. cost derived as Full cap. − Base (no interception-only measurement for Triton; includes
post-capture Python overhead).
‡ vLLM: Full cap. < Base because kerncap extract terminates the host (∼164 s inference
benchmark) once the target dispatch is captured (∼61 s). Cap. cost estimated as 30 GB ÷ 1.7 GB/s ≈ 18 s using the HIP snapshot bandwidth from this table.
Cross-architecture overhead. Table 4 summarizes the same overhead decomposition for the three HIP workloads on each architecture. Interception-only overhead is consistently ≤ 1.2× on the two CDNA generations and ≤ 1.4× on RDNA3 for the larger workloads, with the rocBLAS micro-benchmark’s elevated ratio (1.34– 2.04×) reflecting its extremely short baseline (3.9–7.2 s) rather than a higher absolute per-dispatch cost. Capture cost is bandwidthbound by hsa_memory_copy: on LAMMPS it falls within ∼1.5 s across all three archs for an ∼8.4 GB snapshot (∼1.3–1.7 GB/s effective bandwidth), while llama.cpp shows materially lower snapshot throughput on gfx1100 (∼830 MB/s vs. ∼1.7 GB/s on gfx942), which we attribute to the consumer-class Radeon’s PCIe topology rather than to the capture path itself. Most importantly, the capture mechanism completes correctly and within seconds on every workload-architecture pair we tested. Table 4: Cross-architecture overhead for the HIP workloads (𝑁 =10, 1 warmup per cell). Snapshot sizes from Table 2. Workload
Arch
Base (s)
Inter.
Full
Cap. (s)
rocprofv3
llama.cpp
gfx942 gfx90a gfx1100
8.30 8.99 10.27
1.06× 1.15× 1.18×
1.90× 2.59× 2.53×
7.0 13.0 13.8
1.45× 1.64× 1.60×
LAMMPS
gfx942 gfx90a gfx1100
9.96 21.59 44.05
1.04× 1.02× 1.01×
1.54× 1.32× 1.16×
4.9 6.4 6.4
1.12× 1.06× 1.06×
rocBLAS GEMM
gfx942 gfx90a gfx1100
3.90 7.18 5.49
2.04× 1.75× 1.34×
2.10× 1.79× 1.37×
0.2 0.3 0.2
1.72× 1.62× 1.52×
Inter. = interception-only overhead; Full = instrumented full-app overhead; Cap. = capture cost (full − interception, in seconds); rocp. = rocprofv3 --kernel-trace overhead for comparison.
Wall-clock overhead in context. Capture cost is fixed per-invocation rather than proportional to application duration, so a 10-second HIP baseline sees its multi-GB snapshot dominate total runtime (1.5– 1.9×, Table 3), while vLLM’s ∼2.7-minute run amortizes the same
Ramos, and Lowery
class of fixed cost (30 GB ÷ 1.7 GB/s ≈ 18 s, against a 164 s baseline) to roughly 11% of total runtime—several times smaller than the 1.5– 1.9× ratios on the short-baseline workloads. The elevated rocBLAS interception ratio (2.04×) illustrates that the per-dispatch cost is absolute, not proportional: workloads with many short kernels show worse ratios at identical per-call cost, so for production workloads running minutes to hours the interception tax is negligible in practice. To isolate steady-state compute throughput from process initialization, we extract llama-bench’s self-reported token-generation rate (tg32) across all configurations (Table 5). Prompt-processing (pp512) is stable within 0.7% across all four settings and is omitted from the table. Table 5: llama.cpp token-generation throughput (median ± std. dev., 𝑁 =10). rocprofv3 --kernel-trace shown for comparison. Configuration Baseline Interception only Full capture rocprofv3 --kernel-trace
tg32 (tokens/s) 254 ± 7 251 ± 29 195 ± 35 158 ± 42
Table 5), because per-dispatch timestamp collection introduces serialization that disproportionately affects short, latency-sensitive kernels. Kerncap’s interception, by contrast, preserves tg32 throughput while providing full kernel-extraction capability.
5.4
Optimization Workflow
The central claim of this paper is that Kerncap transforms GPU kernel optimization from a slow, full-application loop into a fast, isolated edit-recompile-validate loop. We quantify this by comparing the two workflows on two workloads (llama.cpp and LAMMPS) on gfx942, focusing on the inner edit-build-replay loop that developers actually iterate on between edits. Validation is reported separately (below) to avoid conflating routine functional checks with heavyweight hardware-counter profiling. Table 6: Inner-loop iteration time comparison: traditional full-application workflow vs. Kerncap isolated reproducer workflow, for two HIP workloads on gfx942. “Build” is a source-level edit’s incremental rebuild; “Run” is the wallclock cost of exercising the kernel after the rebuild. Validation is excluded from the inner loop and discussed separately. llama.cpp Step
Interception-only preserves tg32 within noise (251 vs. baseline 254), confirming that the per-dispatch callback mechanism does not degrade steady-state compute. Full capture drops tg32 to 195 (−23%), with an elevated standard deviation (35 vs. baseline 7) indicating a transient one-time snapshot cost rather than a sustained degradation—the developer captures once and iterates on the isolated reproducer thereafter. rocprofv3 imposes a larger sustained penalty (−38%) on the same metric because its per-dispatch timestamping serializes short, latency-sensitive kernels. Capture cost breakdown. On gfx942, HIP capture cost tracks device-memory throughput at ∼1.7 GB/s (combined hsa_memory_copy and disk writeback) and is independent of kernel count or application structure. Cross-architecture behavior (Table 4) is consistent except on the consumer-class W7900, where llama.cpp’s snapshot slows to ∼830 MB/s—attributable to that platform’s PCIe topology, not to the capture path. Triton extraction times (Table 3) bundle post-capture Python overhead and short-circuit teardown, so their raw I/O throughput is not directly comparable. Comparison with rocprofv3. On gfx942, rocprofv3 --kernel-trace adds 1.45× wall-clock overhead on llama.cpp and 1.12× on LAMMPS—roughly comparable to Kerncap’s interception-only overhead (1.04–1.06×). On vLLM, rocprofv3 imposes 1.33× overhead, while Kerncap’s short-circuit extraction completes in 60.8 s (well under baseline) because per-dispatch timestamp recording scales with kernel count (3.1 M dispatches for vLLM vs. 110 K for llama.cpp). The same pattern holds across architectures (Table 4): rocprofv3 consistently imposes higher per-dispatch overhead than Kerncap’s interception path for the larger workloads. More importantly, rocprofv3 significantly degrades token-generation throughput on llama.cpp (−38% for tg32,
LAMMPS
Traditional Kerncap Traditional Kerncap
Edit source Build Run / replay
— 128 s 8.3 s
— 18.3 s 6.9 s
— 64 s† 10.0 s
— 13.8 s 4.3 s
Inner-loop total Speedup
∼136 s
∼25 s
∼74 s
∼18 s
5.4×
4.1×
† LAMMPS Kokkos rebuild time, single-file edit (measured on the same machine).
Validation cost in context. The inner-loop comparison above intentionally excludes validation, because the validation step a developer runs between edits is rarely the same step they run before promoting an optimization to production. Kerncap supports both points on this spectrum. A lightweight smoke-test validation (kerncap validate default mode) confirms that the captured kernel replays without error in 4.6 s on llama.cpp and 2.7 s on LAMMPS—suitable for routine inter-edit checks. A strict byte-exact validation, comparing every output region against the captured baseline, takes 129.4 s on llama.cpp—suitable for final correctness sign-off. By comparison, the traditional workflow’s equivalent step— running the full application under rocprof-compute to collect hardware counters—takes 2,072 s on llama.cpp due to multiplexing overhead [1], an order of magnitude longer than even the strict Kerncap validation. Across both workloads, Kerncap’s inner-loop speedup is 4–5× for routine iteration and grows to ∼14× when end-to-end correctness validation is included. Case study: optimizing a llama.cpp kernel. We demonstrate the full Kerncap workflow on the mul_mat_vec_q kernel, which accounts for 8.0% of total GPU time in llama.cpp inference (7,740 calls, as reported by kerncap profile). (1) Extract. kerncap extract mul_mat_vec_q --cmd "./llama-bench -m gpt-oss-20b-mxfp4.gguf -p 512 -n 32 -ngl 99 -fa 1" --source-dir ./llama_cpp
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
completes in 15.6 seconds on gfx942, producing a reproducer with 6 source files, 21 memory regions (12,233 MB snapshot), and a Makefile with VFS overlay. (2) Baseline replay. kerncap replay ./isolated/mul_mat_vec_q --iterations 10 reports an average kernel time of 34.5 𝜇s. (3) Edit. We refactor the main block-accumulation loop in kernel_variant.cpp by hoisting a runtime branch (if (use_gate)) completely outside the tightly unrolled inner loop and forcing the compiler to inline the dot-product function via a templated if constexpr evaluation. This eliminates branch overhead and prevents register spilling, allowing the compiler to generate a perfectly unrolled, branch-free inner loop. (4) Recompile. make recompile rebuilds the single HSACO in 18.3 seconds (vs. 128 seconds for a full llama.cpp CMake rebuild). (5) Variant replay. kerncap replay ./isolated/mul_mat_vec_q --hsaco optimized.hsaco --iterations 10 reports 28.3 𝜇s. Compared to the baseline 34.5 𝜇s, this represents a 1.22× speedup (an 18% reduction in execution time). The ability to implement, test, and verify this hypothesis in seconds demonstrates the value of the isolated iteration loop. (6) Validate. kerncap validate ./isolated/mul_mat_vec_q --hsaco optimized.hsaco confirms byte-exact output match across all 21 compared memory regions in 129.4 s. As discussed above, the equivalent step in the traditional workflow—running the full application under ROCm Compute Profiler to collect hardware counters, which takes 2,072 seconds due to multiplexing overhead [1]—is an order of magnitude slower. The end-to-end cycle—edit, recompile, replay, validate—completes in ∼162 seconds on gfx942, compared to ∼2,208 seconds (nearly 37 minutes) for the equivalent traditional workflow (a 13.6× speedup). When the heavyweight final-validation step is omitted in favor of the smoke-test mode used during rapid iteration, the inner-loop speedup is 5.4× for llama.cpp (Table 6) and 4.1× for LAMMPS. Iteration count in context. The final 1.22× speedup above was reached after five edit-recompile-validate iterations. Under the traditional workflow, five iterations of the inner loop cost 5 × 136 ≈ 680 s (∼11 min) of rebuild and application-launch time; under Kerncap, the same five iterations complete in 5 × 25 ≈ 125 s (∼2 min), a direct savings of ∼9 minutes on this single kernel. Including the heavyweight byte-exact validation step above, the full optimization campaign costs ∼254 s (∼4 min) under Kerncap versus ∼2,752 s (∼46 min) traditionally—a ∼42 minute savings per kernel. Scaled across even a modest optimization sweep of ten kernels at five iterations each, Kerncap reclaims roughly seven hours of developer wall time, converting kernel optimization from an overnight task into an interactive one.
6
Related Work
Kernel-level replay and checkpointing. The closest prior art is NVIDIA’s kernel-replay infrastructure. Nsight Compute [9] supports kernel-replay, range-replay, and application-replay modes that re-execute regions of GPU code to multiplex hardware-counter collection across passes. CUPTI exposes the same replay machinery to third-party tools, with the CUDA-11.5+ Checkpoint API [8] explicitly designed to save and restore device state across replay passes; this is the direct conceptual analog of Kerncap’s device-memory snapshot. Three differences separate them. (1) Nsight Compute and CUPTI are closed-source NVIDIA-only tools; Kerncap is open and AMD-native. (2) Their captured state is consumed internally for counter replay, not exposed as an editable, rebuildable artifact for the developer. (3) The Checkpoint API restores state via explicit per-allocation save/restore hooks; Kerncap’s address-space closure (Section 4.2) captures the whole address space at once, sidestepping the need to enumerate individual allocations or chase embedded device pointers. Interactive GPU debuggers. CUDA-GDB and ROCm’s rocgdb provide interactive debugging of GPU kernels but focus on singlestepping and breakpoints rather than kernel isolation and reproduction [7]. NVIDIA’s Compute Sanitizer detects memory errors but does not capture state for replay. These tools complement Kerncap by operating at different points in the optimization workflow. Compiler-emitted reproducer artifacts. PyTorch’s TORCH_COMPILE_DEBUG flag emits the generated Triton/Inductor source for compiled regions, including a self-contained output_code.py that can be re-executed standalone [12]. This is, in effect, a de-facto Triton reproducer for torch.compile-generated kernels and overlaps with Kerncap’s Triton path. However, it is restricted to kernels that Inductor itself generates—not arbitrary hand-written Triton kernels such as Flash Attention—and it does not capture runtime tensor inputs (only shapes and dtypes); any validation it offers is a best-effort eager-vs-compiled allclose comparison rather than a replay against captured runtime state. Kernel benchmarking frameworks. Triton’s built-in benchmarking utilities and AMD’s hipBench provide performance measurement infrastructure but require the developer to manually construct the kernel launch harness, allocate buffers, and populate inputs. Kerncap automates this construction. HSA interception tools. AMD’s rocprofiler-sdk [2] and roctracer [3] provide HSA API table interception for performance counter collection and API tracing, respectively. Kerncap leverages the same rocprofiler-sdk intercept table registration framework to obtain the HSA API table, but applies it to a fundamentally different problem: complete kernel state capture—including arguments, device memory, and source code—for reproducer generation. Record-and-replay systems. Process-level record-and-replay tools such as rr and CRIU [10, 14] capture entire process state (CPU registers, memory maps, file descriptors) for deterministic replay. However, they operate at the CPU/OS level and do not capture GPU state (device memory, kernel dispatch parameters, HSA queues). Kerncap provides GPU-kernel-level replay, capturing only the state needed for a single kernel dispatch.
Ramos, and Lowery
Build system and compilation database tools. Kerncap’s source discovery leverages compile_commands.json, the compilation database format standardized by Clang [15]. Kerncap extends this information with runtime dispatch data to resolve template instantiation ambiguities that static analysis alone cannot resolve. AMD GPU optimization ecosystem. Several AMD-adjacent tools occupy related niches in the GPU performance-engineering workflow. ROCm Compute Profiler (rocprof-compute) collects hardware counters at application granularity [1]; we use it as the validation baseline in our case study (Section 5.4). Recent agent-based kernel generation systems—such as GEAK [18] and the KernelBench [11] benchmark—require a fast inner loop that evaluates candidate kernels in isolation against fixed, realistic inputs. Kerncap produces exactly this kind of self-contained, validated reproducer, positioning the tool as infrastructure for both human optimization workflows and automated kernel-generation agents. Summary. The combinatorial position Kerncap occupies is, to our knowledge, unoccupied by any existing open tool: HIP and Triton interception, VA-faithful device-memory replay, and autotunerpinned validation are each available in isolation, but no single artifact bundles all four. This is precisely the substrate that an automated kernel-generation loop—human or agentic—requires: a fast, isolated, validated evaluation harness for arbitrary candidate kernels.
7
Limitations and Future Work
Single-kernel isolation. Kerncap captures one kernel dispatch at a time. It does not track inter-kernel data dependencies or stream synchronization, so kernels that depend on the output of preceding kernels require those predecessors to have already executed before the capture point. In practice, this is usually the case because Kerncap intercepts during normal application execution, but capturing a sequence of dependent kernels for joint replay remains future work. Single-GPU.. The current implementation targets single-GPU applications. Multi-GPU dispatch capture would require coordinating libkerncap.so instances across multiple HSA agents and merging their memory snapshots. Memory snapshot fragility. If the application frees a device buffer between kernel completion and the memory snapshot, hsa_memory_copy may fail for that region. Kerncap mitigates this by writing metadata files before the snapshot and tolerating perregion copy failures, but the affected region’s data will be lost. No host-side state capture. Kerncap captures GPU-side state (device memory, kernarg buffers, code objects) but not host-side state (CPU memory, file handles, environment variables). Kernels that read from host-mapped memory may not replay correctly. Source discovery limitations. HIP source discovery works best when compile_commands.json is available; without it, the finder falls back to heuristics that may miss complex build configurations. The #include tracer handles local includes but not computed includes or macro-generated include paths.
GPU portability layers. Portability-layer kernels (Kokkos, RAJA, SYCL [13]) are captured and replayed correctly via DWARF-based source discovery (Section 4.3) and module-variable restoration (Section 4.2), but the reproducer’s editable source is the framework template expansion rather than the user-level functor or lambda the developer authored. Tracing through framework abstractions back to application-level code remains future work. Hardware-specific autotuner configurations. Triton autotuner configurations are hardware-specific. A configuration captured on MI300X may not be valid on MI250X due to differences in shared memory size, wavefront width, or register pressure. Cross-GPU reproducer portability requires re-autotuning or manual config adjustment. Limited RDNA evaluation. Our RDNA evaluation is scoped to the three HIP workloads (llama.cpp, LAMMPS, rocBLAS GEMM) on a single gfx1100 system; with more time, we would have liked to have performed a full benchmark campaign across the Triton workloads or newer RDNA hardware. More rigorous RDNA characterization— particularly on the RDNA 4 (Navi 4, gfx12xx) family—remains future work. Future directions. We plan to extend Kerncap in several directions: (1) multi-kernel capture for dependent kernel sequences with automatic dependency tracking; (2) integration with kernel autotuning tools for closed-loop optimization; (3) support for NVIDIA CUDA via analogous CUPTI or driver API interception; (4) a persistent kernel database for regression testing across ROCm versions.
8
Conclusion
We presented Kerncap, a tool for automated GPU kernel extraction and isolation on AMD GPUs. Kerncap addresses the labor-intensive manual process of isolating individual kernels from complex GPU applications by automating five key tasks: runtime interception at the HSA level (unified for both HIP and Triton, with a Python compile-hook shim bridging Triton-level metadata to the HSAlevel capture), VA-faithful device memory snapshot that inherently preserves embedded pointer relationships, automated source discovery with translation unit resolution and dependency tracing, self-contained reproducer generation with VFS overlay recompilation support, and correctness validation. Together, these five tasks recover the three reproducer components introduced in Section 1 (Figure 1): the kernel definition (HSACO capture, source discovery, and reproducer assembly), the runtime state (the VA-faithful devicememory snapshot), and the environment (VFS-overlay recompile and HSA replay). The tool unifies HIP and Triton capture at the HSA layer while accommodating their fundamentally different kernel-dispatch models (ahead-of-time compiled HSACOs vs. Python-JIT’d code objects) through a lightweight Python metadata shim, presenting a single kerncap extract CLI to the developer. Two principles underpin the design. Address-space closure—an address space is a closure over its pointer graph—eliminates the need for DWARF metadata or pointer analysis, handling arbitrarily complex argument layouts including double-pointer indirection. Tuning-pinned reproducers preserve the implicit numerical contract that a JIT-compiled Triton kernel carries between its autotuner-selected tile sizes and its
Kerncap: Automated Kernel Extraction and Isolation for AMD GPUs
outputs. These principles are realized through a Clang VFS overlay that enables source-level kernel editing and recompilation using the exact original build flags without any build system modifications. Our case study on llama.cpp demonstrates the practical impact: the isolated edit-recompile-validate loop completes in 162 seconds compared to ∼37 minutes for the traditional full-application workflow—a 13.6× speedup. By reducing kernel isolation from hours of manual effort to a single command, Kerncap enables faster GPU kernel optimization cycles and lowers the barrier to kernel-level performance engineering on AMD hardware. More broadly, the artifact Kerncap produces—a self-contained, validated, editable reproducer—is precisely the substrate that both human experts and emerging LLM-driven kernel-generation agents need to iterate on GPU code in isolation.
Acknowledgments This work was supported by Advanced Micro Devices, Inc. under the AMD AI & HPC Cluster Program. The authors would like to thank Karl Schulz, Mehdi Saeedi, Madhu Srinivasan, and Ralph Wittig for their support and guidance. AMD, the AMD Arrow logo, AMD CDNA, AMD Instinct, AMD ROCm, AMD Infinity Cache, AMD Infinity Fabric, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
References [1] AMD. 2026. ROCm Compute Profiler Documentation. https://rocm.docs.amd. com/projects/rocprofiler-compute/en/latest/ ROCm Compute Profiler is a kernellevel profiling tool for machine learning and high performance computing (HPC) workloads running on AMD Instinct™ accelerators. [2] AMD. 2026. ROCProfiler-SDK Documentation. https://rocm.docs.amd.com/ projects/rocprofiler-sdk/en/latest/ ROCprofiler-SDK is a tooling infrastructure for profiling general-purpose GPU compute applications running on the ROCm software. [3] AMD. 2026. ROCTracer Documentation. https://rocm.docs.amd.com/projects/ roctracer/en/latest/ ROCTracer consists of the ROCTracer and ROC-TX libraries, which provide APIs to help you trace an application in the runtime. [4] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. 2022. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. arXiv:2205.14135 [cs.LG] https://arxiv.org/abs/2205.14135 [5] Georgi Gerganov and contributors. 2023. llama.cpp: LLM inference in C/C++. https://github.com/ggml-org/llama.cpp. Accessed: 2026-04-17. [6] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph E. Gonzalez, Hao Zhang, and Ion Stoica. 2023. Efficient Memory Management for Large Language Model Serving with PagedAttention. In Proceedings of the ACM SIGOPS 29th Symposium on Operating Systems Principles. [7] NVIDIA. 2026. CUDA-GDB: The NVIDIA CUDA Debugger. https://developer. nvidia.com/cuda-gdb NVIDIA Developer Documentation. [8] NVIDIA. 2026. CUPTI: Checkpoint API. https://docs.nvidia.com/cupti/api/ group__CUPTI__CHECKPOINT__API.html NVIDIA CUPTI Checkpoint API Documentation. [9] NVIDIA. 2026. NVIDIA Nsight Compute: Kernel Profiling Guide. https://docs. nvidia.com/nsight-compute/index.html NVIDIA Developer Documentation. [10] Robert O’Callahan, Chris Jones, Nathan Froyd, Kyle Huey, Albert Noll, and Nimrod Partush. 2017. Engineering Record And Replay For Deployability: Extended Technical Report. arXiv:1705.05937 [cs.PL] https://arxiv.org/abs/1705.05937 [11] Anne Ouyang, Simon Guo, Simran Arora, Alex L. Zhang, William Hu, Christopher Ré, and Azalia Mirhoseini. 2025. KernelBench: Can LLMs Write Efficient GPU Kernels? arXiv:2502.10517 [cs.LG] https://arxiv.org/abs/2502.10517 [12] Adam Paszke, Sam Gross, Francisco Massa, Adam Lerer, James Bradbury, Gregory Chanan, Trevor Killeen, Zeming Lin, Natalia Gimelshein, Luca Antiga, Alban Desmaison, Andreas Köpf, Edward Yang, Zach DeVito, Martin Raison, Alykhan Tejani, Sasank Chilamkurthy, Benoit Steiner, Lu Fang, Junjie Bai, and Soumith Chintala. 2019. PyTorch: an imperative style, high-performance deep learning library. Curran Associates Inc., Red Hook, NY, USA.
[13] Sivasankaran Rajamanickam, Seher Acer, Luc Berger-Vergiat, Vinh Dang, Nathan Ellingwood, Evan Harvey, Brian Kelley, Christian R. Trott, Jeremiah Wilke, and Ichitaro Yamazaki. 2021. Kokkos Kernels: Performance Portable Sparse/Dense Linear Algebra and Graph Kernels. arXiv:2103.11991 [cs.MS] https://arxiv.org/ abs/2103.11991 [14] Radostin Stoyanov, Viktória Spišaková, Jesus Ramos, Steven Gurfinkel, Andrei Vagin, Adrian Reber, Wesley Armour, and Rodrigo Bruno. 2025. CRIUgpu: Transparent Checkpointing of GPU-Accelerated Workloads. arXiv:2502.16631 [cs.DC] https://arxiv.org/abs/2502.16631 [15] The Clang Team. 2026. JSON Compilation Database Format Specification. https: //clang.llvm.org/docs/JSONCompilationDatabase.html Clang Documentation. [16] A. P. Thompson, H. M. Aktulga, R. Berger, D. S. Bolintineanu, W. M. Brown, P. S. Crozier, P. J. in ’t Veld, A. Kohlmeyer, S. G. Moore, T. D. Nguyen, R. Shan, M. J. Stevens, J. Tranchida, C. Trott, and S. J. Plimpton. 2022. LAMMPS - a flexible simulation tool for particle-based materials modeling at the atomic, meso, and continuum scales. Comp. Phys. Comm. 271 (2022), 108171. doi:10.1016/j.cpc.2021. 108171 [17] Philippe Tillet, H. T. Kung, and David Cox. 2019. Triton: an intermediate language and compiler for tiled neural network computations. In Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages (Phoenix, AZ, USA) (MAPL 2019). Association for Computing Machinery, New York, NY, USA, 10–19. doi:10.1145/3315508.3329973 [18] Jianghui Wang, Vinay Joshi, Saptarshi Majumder, Xu Chao, Bin Ding, Ziqiong Liu, Pratik Prabhanjan Brahma, Dong Li, Zicheng Liu, and Emad Barsoum. 2025. Geak: Introducing Triton Kernel AI Agent & Evaluation Benchmarks. arXiv:2507.23194 [cs.CL] https://arxiv.org/abs/2507.23194