Kerncap: Automated GPU Kernel Isolation

• Kerncap is an automated system that isolates GPU kernels by capturing the compiled kernel, runtime state, and build environment to create a reproducible, self-contained reproducer.
• It follows a five-stage extraction pipeline including profiling, state capture, source discovery, reproducer generation, and validation through HSA-level API interception for both HIP and Triton.
• The system significantly speeds up the edit–recompile–validate loop, achieving up to a 13.6× speedup in kernel optimization workflows and supporting robust autotuning and regression testing.

Kerncap is an automated kernel-extraction and isolation system for AMD GPUs that captures a GPU kernel from a larger application and emits a self-contained reproducer preserving the kernel definition, runtime state, and replay environment (Ramos et al., 4 May 2026). It is designed to replace the slow workflow of manually tracing source, reconstructing build flags, recovering dispatch configuration, capturing device-resident state, and rebuilding the full application for each edit with a faster edit–recompile–validate loop. The system operates at the HSA runtime level for both HIP and Triton, uses a virtual-address-faithful snapshot of device memory, and generates replay artifacts that can be recompiled and validated in isolation (Ramos et al., 4 May 2026).

1. Problem formulation and scope

The paper frames kernel isolation as the recovery of three tightly coupled components: the definition, consisting of the compiled kernel binary together with transitive source dependencies and build flags; the runtime state, consisting of dispatch dimensions, kernel arguments, and all device memory the kernel reads; and the environment, consisting of a faithful build and replay setup that preserves the original execution (Ramos et al., 4 May 2026). In this formulation, the difficulty is not limited to locating a kernel symbol. It lies in reconstructing a complete executable context from large codebases, template-heavy build systems, JIT pipelines, and runtime-generated metadata.

Manual isolation is described as brittle because these components are distributed across static and dynamic layers. The developer may need to trace mangled names through headers and templates, reconstruct compile flags and translation units, identify the exact dispatch configuration, capture device-resident state including indirectly referenced memory, and rebuild the entire application merely to test a small kernel edit (Ramos et al., 4 May 2026). The llama.cpp example is used to make the bottleneck concrete: a full CMake rebuild takes 128 s, so the optimization loop is dominated by rebuilding rather than by the kernel modification itself (Ramos et al., 4 May 2026).

A central point in the paper is that kernel optimization is bottlenecked less by writing the optimization itself than by isolating the kernel from the host application. This suggests that the contribution of Kerncap is methodological as much as infrastructural: it redefines kernel tuning as a reproducibility problem at the level of dispatch state, memory image, and source provenance rather than as a purely compilation-centric task.

2. End-to-end extraction pipeline

Kerncap is exposed through a CLI with three commands:

kerncap profile kerncap extract kerncap replay

The main work occurs in extract, within a five-stage pipeline (Ramos et al., 4 May 2026). The first stage, Profile, uses rocprofv3 --kernel-trace to identify target kernels, although this is optional when the target is already known. The second stage, Capture runtime state, loads libkerncap.so via LD_PRELOAD and interposes on AMD’s HSA runtime using ROCProfiler-SDK’s HSA table interception, which supports both HIP and Triton. The third stage, Discover source, uses different mechanisms for HIP and Triton. For HIP, Kerncap consults compile_commands.json, DWARF line information when available, and fallback grep/include tracing to recover editable source and the correct translation unit. For Triton, it parses Python ASTs and import graphs to locate the @triton.jit or @triton.autotune kernel together with its dependencies (Ramos et al., 4 May 2026).

The fourth stage, Generate a reproducer, emits a self-contained project. HIP reproducers contain a Clang VFS overlay and a Makefile; Triton reproducers are standalone Python replay scripts with tuning pinned (Ramos et al., 4 May 2026). The fifth stage, Validate / replay, replays the kernel through raw HSA, restores the device address space, and checks outputs. HIP uses either a smoke test or a byte-exact diff, whereas Triton uses numpy.allclose with NaN detection (Ramos et al., 4 May 2026).

This organization matters because it makes explicit that extraction is not only about code generation. The replay artifact must encode dispatch metadata, memory layout, code objects, and validation logic. A plausible implication is that Kerncap’s real unit of capture is the executed kernel instance rather than the source file alone.

3. Runtime interception and metadata recovery

HIP path

For HIP, Kerncap hooks the HSA API table at runtime through ROCProfiler-SDK. libkerncap.so registers via rocprofiler_at_intercept_table_registration, receives the live HsaApiTable, and replaces selected function pointers with wrappers while forwarding calls onward (Ramos et al., 4 May 2026). It intercepts queue creation through hsa_queue_create in order to install per-packet callbacks via hsa_amd_queue_intercept_create and hsa_amd_queue_intercept_register. It also intercepts memory allocation APIs to track device allocations and their sizes, symbol queries to map kernel object handles to mangled names, and code-object loading to capture HSACO blobs (Ramos et al., 4 May 2026). Kernel dispatch metadata is obtained directly from the HSA packet, including grid size, workgroup size, kernarg pointer, and symbol identity.

Triton path

Triton kernels also dispatch through HSA, so the same HSA-level interception captures their execution (Ramos et al., 4 May 2026). The missing information is Python-side metadata that HSA does not expose: the user-visible kernel name, typed Python signature, tl.constexpr bindings, autotuner-selected configuration, and tensor metadata such as dtype and strides. Kerncap bridges this gap with a lightweight Python compile-hook shim: a sitecustomize.py hook installs itself at interpreter startup, intercepts Triton JIT compilation, records metadata into name_map.json, and keys the record by the SHA-256 of the emitted HSACO (Ramos et al., 4 May 2026).

A C++ kernarg metadata parser then reads AMDGPU code-object metadata and matches binary kernarg slots to typed arguments, offsets, sizes, and constexpr values, allowing the HSA dispatch to be reassembled into the original Triton call (Ramos et al., 4 May 2026). The paper is explicit that HSA state alone is insufficient for reconstructing a Triton replay artifact. The shim supplies the high-level semantics of the captured kernarg bytes.

This division between low-level capture and high-level metadata recovery is important for interpreting the system. A common misconception is that JIT kernels can be reproduced from a captured binary alone. The paper rejects that view for Triton: the original call contract depends on metadata and autotuner state that are external to the HSA dispatch record.

4. Address-space closure and replay invariants

A defining concept in Kerncap is address-space closure, implemented through a VA-faithful capture of device memory (Ramos et al., 4 May 2026). Kerncap snapshots every tracked device allocation at its original virtual address. The paper states the principle as follows: “an address space is a closure over its pointer graph” (Ramos et al., 4 May 2026). The operational consequence is that if one allocation contains a pointer to another, and that second allocation contains a pointer to a third, restoring all of them at the same virtual addresses preserves the embedded device-pointer graph without DWARF metadata, recursive pointer chasing, or special handling for nested device pointers such as T** (Ramos et al., 4 May 2026).

Memory capture proceeds by iterating over the tracked allocation map, streaming each region to disk in chunks via hsa_memory_copy, storing metadata in memory_regions.json, and storing the kernarg buffer separately in kernarg.bin (Ramos et al., 4 May 2026). Kerncap also captures HSA module variables, including __constant__ memory and similar executable-load-time variables, restoring them before replay dispatch (Ramos et al., 4 May 2026). The paper identifies this mechanism as the enabler for the vLLM Mixture-of-Experts case, in which the kernel dereferences multiple levels of indirection through expert weight tables and the capture therefore includes the large weight pool rather than only immediate arguments (Ramos et al., 4 May 2026).

Replay is formalized as a VA-faithful HSA-level process: parse memory_regions.json and dispatch.json; pre-reserve virtual address ranges with mmap(MAP_FIXED_NOREPLACE); initialize HSA; allocate and map memory at the exact captured addresses using HSA VMEM APIs; restore region contents; load HSACO, resolve the symbol, fill kernarg, and dispatch; optionally dump outputs for validation (Ramos et al., 4 May 2026). The critical invariant is that replay must use the same virtual addresses. If allocation returns a different VA, Kerncap aborts (Ramos et al., 4 May 2026).

The paper also delineates the limits of closure. Host-resident pointers in kernel arguments are not covered, host ASLR can break host-side pointers passed in kernargs, and pointers the runtime never tracked are outside the closure (Ramos et al., 4 May 2026). This is not presented as a flaw in the mechanism; rather, it defines the boundary of what a device-side VA-faithful replay can guarantee.

5. Reproducer generation for HIP and Triton

HIP and Triton require different reproducer strategies because the relation between source, build state, and runtime state is different in the two ecosystems (Ramos et al., 4 May 2026).

For HIP, Kerncap generates a self-contained project in which capture/ contains dispatch metadata, HSACO, kernarg, and all memory regions, while source files and headers are copied locally (Ramos et al., 4 May 2026). A vfs.yaml overlay maps those local copies back to the original filesystem paths. The paper identifies the Clang Virtual File System overlay as the key mechanism for source-level recompilation without changing the original build system (Ramos et al., 4 May 2026). The generated Makefile reuses the exact original compile command from compile_commands.json, augmented by:

-ivfsoverlay vfs.yaml --cuda-device-only --no-gpu-bundle-output

This preserves the original include paths, optimization flags, and architecture settings while allowing edits to the copied sources in isolation (Ramos et al., 4 May 2026).

For Triton, the central issue is not only source recovery but preservation of the autotuner-selected configuration (Ramos et al., 4 May 2026). The paper argues that Triton kernels have an implicit numerical contract between tuning state and output. When the original kernel used @triton.autotune, the reproducer therefore bypasses autotuning and calls the kernel directly with the captured winning configuration (Ramos et al., 4 May 2026). This is described as a tuning-pinned reproducer.

The FlashAttention example makes the motivation concrete: switching to the second-fastest configuration changes 11.3% of output elements, even though it is only 7.7% slower (Ramos et al., 4 May 2026). The paper’s conclusion is that a faithful Triton reproducer must bind the captured autotuner configuration into the artifact in order to preserve the JIT kernel’s numerical contract.

This distinction between HIP and Triton shows that Kerncap is not a uniform dump-and-replay system. Its artifacts are specialized to the semantics of the captured kernel stack: source-level recompilation for HIP and tuning-pinned replay for Triton.

6. Experimental evaluation and performance characteristics

The evaluation covers six workloads across three AMD GPU architectures: CDNA3 (MI300X / gfx942), CDNA2 (MI210 / gfx90a), and RDNA3 (W7900 / gfx1100) (Ramos et al., 4 May 2026). The workloads are llama.cpp (HIP), LAMMPS (HIP), rocBLAS GEMM (HIP, no source recovery needed), Flash Attention 2 (Triton), vLLM (Triton), and torch.compile (Triton / Inductor-generated) (Ramos et al., 4 May 2026).

Workload	Stack	Snapshot size
llama.cpp	HIP	about 11.5–12.2 GB
LAMMPS	HIP	about 8.4 GB
rocBLAS GEMM	HIP	152–248 MB
Flash Attention 2	Triton	178 MB
vLLM	Triton	about 29.3–30.1 GB
torch.compile	Triton / Inductor-generated	227 MB

Across these experiments, snapshot sizes ranged from 152 MB to 30 GB (Ramos et al., 4 May 2026). The vLLM case is especially notable: Kerncap captures 185–186 memory regions totaling about 30 GB, including the Mixture-of-Experts weight pool reached through pointer indirection (Ramos et al., 4 May 2026). The paper reports successful extraction and validation on every workload and architecture tested, with minor cross-architecture variation in region count attributed to allocation behavior rather than capture errors (Ramos et al., 4 May 2026).

Validation behavior differs by stack but is uniformly successful in the reported experiments. HIP workloads replay byte-identically, Triton workloads validate with zero error under allclose, and the system preserves wavefront-width differences between CDNA and RDNA (Ramos et al., 4 May 2026). Interception-only overhead is reported as low, generally near 1.0–1.2× for larger workloads, and the paper notes that much of the apparent wall-clock cost comes from one-time snapshotting rather than sustained slowdown (Ramos et al., 4 May 2026). It also reports that rocprofv3 --kernel-trace imposes more sustained overhead than Kerncap’s interception path (Ramos et al., 4 May 2026).

The main case study concerns llama.cpp’s mul_mat_vec_q kernel. The traditional workflow takes about 2,208 s, whereas the Kerncap workflow takes about 162 s, yielding a 13.6× speedup (Ramos et al., 4 May 2026). Within that workflow, make recompile takes 18.3 s compared with 128 s for a full llama.cpp rebuild (Ramos et al., 4 May 2026). The optimized kernel improves from 34.5 μs to 28.3 μs, a 1.22× speedup for the kernel itself (Ramos et al., 4 May 2026). The paper interprets these numbers as evidence that reducing isolation cost can materially accelerate iterative optimization.

7. Interpretation, limitations, and downstream uses

The paper’s broader claim is that the generated reproducers are not only debugging artifacts but a general substrate for automated optimization (Ramos et al., 4 May 2026). Because they provide a self-contained, validated execution harness with fixed runtime inputs, preserved memory state, dispatch configuration, and deterministic replay behavior, they are positioned as useful for autotuning agents, LLM-driven kernel generation workflows, regression testing across compiler or ROCm versions, and hardware-counter or performance studies without rerunning the entire application (Ramos et al., 4 May 2026).

This suggests that Kerncap reifies the kernel instance as a compact benchmark object. In that interpretation, capture is a way of transforming a kernel embedded in a large heterogeneous program into a stable optimization target with preserved semantics. The paper’s emphasis on byte-identical HIP replay, allclose-validated Triton replay, and tuning-pinned JIT artifacts supports this reading (Ramos et al., 4 May 2026).

The limitations are also explicit. Address-space closure does not include host-resident pointers in kernel arguments, can be disrupted by host ASLR for host-side pointers passed in kernargs, and excludes pointers the runtime never tracked (Ramos et al., 4 May 2026). Triton reconstruction further depends on the Python compile-hook shim because HSA does not expose high-level Python signature information, tl.constexpr bindings, or autotuner state (Ramos et al., 4 May 2026). These constraints indicate that Kerncap’s guarantees are strongest for device-resident state observed through the HSA runtime and for JIT environments whose metadata can be intercepted during compilation.

Within those boundaries, the paper presents Kerncap as an automated answer to a precise systems problem: faithful extraction of GPU kernels from large HIP and Triton applications into isolated, editable, and validated reproducer projects (Ramos et al., 4 May 2026).