---
title: 'Kerncap: Automated GPU Kernel Isolation'
url: https://www.emergentmind.com/topics/kerncap
type: topic
---

# Kerncap: Automated GPU Kernel Isolation

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 [2605.03208]. 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 [2605.03208].

## 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 [2605.03208]. 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 [2605.03208]. 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 [2605.03208].

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:

```bash
kerncap profile
kerncap extract
kerncap replay
```

The main work occurs in `extract`, within a five-stage pipeline [2605.03208]. 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 [2605.03208].

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 [2605.03208]. 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 [2605.03208].

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 [2605.03208]. 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 [2605.03208]. 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 [2605.03208]. 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 [2605.03208].

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 [2605.03208]. 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 [2605.03208]. 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”** [2605.03208]. 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**` [2605.03208].

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` [2605.03208]. Kerncap also captures HSA module variables, including `__constant__` memory and similar executable-load-time variables, restoring them before replay dispatch [2605.03208]. 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 [2605.03208].

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 [2605.03208]. The critical invariant is that replay must use the same virtual addresses. If allocation returns a different VA, Kerncap aborts [2605.03208].

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 [2605.03208]. 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 [2605.03208].

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 [2605.03208]. 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** [2605.03208]. The generated Makefile reuses the exact original compile command from `compile_commands.json`, augmented by:

```text
-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 [2605.03208].

For **Triton**, the central issue is not only source recovery but preservation of the autotuner-selected configuration [2605.03208]. 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 [2605.03208]. 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** [2605.03208]. 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**) [2605.03208]. 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) [2605.03208].

| 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** [2605.03208]. 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 [2605.03208]. 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 [2605.03208].

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 [2605.03208]. 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 [2605.03208]. It also reports that `rocprofv3 --kernel-trace` imposes more sustained overhead than Kerncap’s interception path [2605.03208].

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** [2605.03208]. Within that workflow, `make recompile` takes **18.3 s** compared with **128 s** for a full llama.cpp rebuild [2605.03208]. The optimized kernel improves from **34.5 μs** to **28.3 μs**, a **1.22× speedup** for the kernel itself [2605.03208]. 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 [2605.03208]. 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 [2605.03208].

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 [2605.03208].

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 [2605.03208]. 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 [2605.03208]. 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 [2605.03208].

Source: https://www.emergentmind.com/topics/kerncap