Skip to content

README.md

Latest commit

 

History

History
executable file
·
176 lines (125 loc) · 6.96 KB

README.md

File metadata and controls

executable file
·
176 lines (125 loc) · 6.96 KB
Raw

kerncap

Kernel extraction and isolation tool for HIP and Triton applications on AMD GPUs.

kerncap profiles a running application, intercepts a target kernel dispatch, captures its complete runtime state (full device memory snapshot, kernarg buffer, HSACO), and generates a standalone reproducer that can replay the kernel in isolation using VA-faithful HSA dispatch.

For full documentation, see amdresearch.github.io/intellikit/tools/kerncap.

How it works

1. Profile          rocprofv3 --kernel-trace --stats → rank kernels by duration
2. Capture          HIP:    LD_PRELOAD=libkerncap.so → intercept target dispatch,
                            snapshot all tracked device memory + kernarg buffer + HSACO
                    Triton: LD_PRELOAD capture with code-object metadata +
                            compile-shim name mapping
3. Find source      HIP: __global__ grep + #include tracing
                    Triton: @triton.jit AST match + import tracing (incl. relative imports)
4. Generate         Jinja2 templates → standalone .hip+Makefile or .py reproducer
5. Validate         Build, run reproducer, np.allclose against captured reference

Install

Builds libkerncap.so from source against the host ROCm 7.0+ (requires hipcc, cmake, HSA headers, rocprofiler-sdk — all present in standard ROCm images). No PyTorch or Triton dependency. No network access needed during the C++ build (nlohmann/json is vendored).

# From local source
pip install .

# Editable install for development
pip install -e .[dev]

Usage

Each operation is available as both a Python API and a CLI command.

Profile

Rank kernels by total GPU execution time.

# Profile and print top kernels
kerncap profile -- ./my_app --args

# Save profile to JSON
kerncap profile --output profile.json -- ./my_app
from kerncap import Kerncap

kc = Kerncap()
profile = kc.profile(["./my_app", "--args"])
for kernel in profile[:5]:
    print(f"{kernel.name}: {kernel.total_duration_ns / 1e6:.1f} ms ({kernel.percentage:.1f}%)")

Extract

Capture a kernel's full runtime state and generate a standalone reproducer.

# HIP with source
kerncap extract mul_mat_q --cmd "..." --source-dir ./ggml/src -D GGML_USE_HIP

# Triton
kerncap extract flash_attn_fwd --cmd "..." --source-dir ./flash_attn

# Capture-only (no source)
kerncap extract mul_mat_q --cmd "..."

# Specific dispatch
kerncap extract gemm_kernel --cmd "..." --dispatch 2
result = kc.extract(
    kernel_name="mul_mat_q",
    cmd=["./llama-bench", "-m", "model.gguf", "-p", "512"],
    source_dir="./ggml/src",
    output="./isolated/mul_mat_q",
    defines=["GGML_USE_HIP", "GGML_CUDA_FA_ALL_QUANTS"],
)

Language detection: kerncap auto-detects whether a kernel is HIP or Triton from --source-dir contents. To override, pass --language hip or --language triton on the CLI (or language="triton" in the Python API).

Triton backend compatibility: The --triton-backend flag is available for compatibility with the legacy Python capture backend. HSA is the new default, and the Python backend will soon be deprecated.

Replay

Replay a captured kernel in isolation.

# Replay with captured HSACO
kerncap replay ./isolated/mul_mat_q

# Replay with a variant HSACO
kerncap replay ./isolated/mul_mat_q --hsaco optimized.hsaco

# Benchmark over multiple iterations
kerncap replay ./isolated/mul_mat_q --iterations 100
# Replay baseline vs variant and compare
baseline = kc.replay("./isolated/mul_mat_q")
variant = kc.replay("./isolated/mul_mat_q", hsaco="optimized.hsaco")
print(f"Speedup: {baseline.timing_us / variant.timing_us:.2f}x")

HIP launch mode: If replay conflicts with rocprofv3 (e.g. when profiling the reproducer itself), pass --hip-launch to use the HIP runtime launch path instead of the default HSA dispatch.

Validate

Check correctness of a reproducer or variant HSACO.

# Smoke test — confirm baseline replays without error
kerncap validate ./isolated/mul_mat_q

# Correctness check — compare variant against captured baseline
kerncap validate ./isolated/mul_mat_q --hsaco optimized.hsaco

# Triton — compare with relaxed tolerance
kerncap validate ./isolated/flash_attn_fwd --tolerance 1e-3 --rtol 1e-2
# Correctness check — compare variant against captured baseline
result = kc.validate("./isolated/mul_mat_q", hsaco="optimized.hsaco")
print("Passed:", result.passed)

Validation modes: For VA-faithful captures, baseline validate is a smoke test only until a rebuilt HSACO is available. Pass hsaco, or let kerncap auto-detect candidate.hsaco / optimized.hsaco, to compare captured vs rebuilt execution byte-for-byte.

kerncap validate <dir> auto-detects rebuilt HSACOs from the edit loop: candidate.hsaco from Triton python3 reproducer.py, or optimized.hsaco from HIP make recompile. In that case validation runs captured vs rebuilt HSACO and reports a byte-exact memory-region summary.

Optimization workflow

When source_dir is provided, extract produces a self-contained project for a tight edit-recompile-validate loop:

cd ./isolated/mul_mat_q

make run            # replay baseline
# edit kernel_variant.cpp and/or deps/
make recompile      # recompile into optimized.hsaco
make run-variant    # replay variant
kerncap validate .   # auto-detects optimized.hsaco for correctness check

See full documentation for the complete Python API workflow and details on the generated project layout.

For Triton captures, edit kernel_variant.py (or the copied source file), rerun python3 reproducer.py to produce candidate.hsaco, then run kerncap validate ..

Notes

  • Source-not-found warnings now distinguish between a wrong --source-dir, a likely Triton kernel routed as HIP, and kernels with no source trail such as Tensile, hand-written assembly, JIT-generated binaries, or vendor HSACO blobs.
  • Captures include module-variable snapshots for constant-memory style launches such as Kokkos hip_parallel_launch_constant_memory, and replay restores only module variables belonging to the captured executable.
  • Large device-memory snapshots stream through a bounded staging buffer. Set KERNCAP_SNAPSHOT_CHUNK_BYTES to tune the chunk size when debugging very large captures.

Project structure

src/kerncap.{hip,hpp}     HSA tool loaded via LD_PRELOAD (rocprofiler-sdk registration)
src/replay.cpp             VA-faithful HSA kernel replay binary (kerncap-replay)
src/kernarg_metadata.*     AMDGPU code-object kernarg metadata parser
kerncap/                   Python package (CLI, profiler, capturer, source finder,
                           reproducer generator, validator)
kerncap/templates/         Jinja2 templates for HIP and Triton reproducers
vendor/                    Vendored nlohmann/json headers
tests/                     Unit + integration tests (see tests/README.md)