Menu
CUDA graph replay benchmark + source trace templates
Install with Codex or Claude Copy this prompt, paste it into Codex, Claude, or another assistant, and let it review the skill page and install it for you.
Based on SOC occupation classification
Profile LLM with ATOM, extract GPU kernels, generate kernel inventory
Extract per-layer kernel sequence from ATOM level=3 traces
Launch ATOM server and run throughput benchmark
Single GPU kernel micro-benchmark with roofline analysis
Deep kernel source-level walkthrough for aiter/atom/CK stack
Map complete kernel pipeline for LLM inference (prefill vs decode)
| name | bench-kernel |
| description | CUDA graph replay benchmark + source trace templates |
| user-invocable | true |
Kernel ๆง่ฝๅๆๅๆบ็ ่ฟฝ่ธช็่ฏฆ็ปๆจกๆฟไธไปฃ็ ็ๆฎตใ
import torch
import triton
NUM_INPUTS = 100
WARMUP = 25
REP = 100
def capture_graph(fn, num_inputs=NUM_INPUTS):
"""Capture CUDA graph over num_inputs calls."""
# Warmup outside graph (trigger JIT / lazy init)
for _ in range(3):
fn()
torch.cuda.synchronize()
stream = torch.cuda.Stream()
with torch.cuda.stream(stream):
graph = torch.cuda.CUDAGraph()
with torch.cuda.graph(graph, stream=stream):
for _ in range(num_inputs):
fn()
torch.cuda.synchronize()
return graph
def bench_graph(graph, num_inputs=NUM_INPUTS):
"""Benchmark graph replay, return per-iteration median time in us."""
def replay():
graph.replay()
t_total = triton.testing.do_bench(replay, warmup=WARMUP, rep=REP, return_mode="median")
return t_total / num_inputs * 1000 # ms -> us per iteration
t = triton.testing.do_bench(fn, warmup=25, rep=100, return_mode="median") * 1000 # us
with torch.profiler.profile(
activities=[torch.profiler.ProfilerActivity.CPU, torch.profiler.ProfilerActivity.CUDA],
record_shapes=True,
) as prof:
fn()
torch.cuda.synchronize()
print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))
# <Kernel Name> Source Trace
## ๆฆ่ฟฐ
ไธๅฅ่ฏ่ฏดๆๆดไฝๆต็จๅๅ
ณ้ฎ shapeใ
## Path 1: <ๅญๆต็จๅ>
- [ๅฝๆฐๅ](aiter/ops/quant.py:258) # ๆณจ้
- [ๅญๅฝๆฐ](aiter/ops/quant.py:413) # ๆณจ้
- [ๆดๆทฑๅฑ](aiter/jit/core.py:968)
- HIP Kernel โ
### HIP/CUDA Kernel: `kernel_name`
- ๆบ็ ไฝ็ฝฎ
- Kernel ๅๆฐ๏ผBlockSize, tile size ็ญ๏ผ
- ๆ ธๅฟๆต็จ๏ผ็ผๅทๅ่กจ๏ผ
## Path 2: <ๅฆไธไธชๅญๆต็จ>
...
## ๅฎๆด็ซฏๅฐ็ซฏๆต็จๅพ
็จ ASCII ๆต็จๅพๅฑ็คบๆฐๆฎๆต่ฝฌใ
a8w8_blockscale_1x128x128_256x16x128x256_16x16_16x16_1x2_...
โ โ โ โ โ โ โ โ โ โ โ โโ NXdlPerWave
โ โ โ โ โ โ โ โ โ โ โโโโโ MXdlPerWave
โ โ โ โ โ โ โ โ โ โโโโโโโโ NPerXDL
โ โ โ โ โ โ โ โ โโโโโโโโโโโโโ MPerXDL
โ โ โ โ โ โ โ โโโโโโโโโโโโโโ KPerBlock
โ โ โ โ โ โ โโโโโโโโโโโโโโโโโโโโ NPerBlock
โ โ โ โ โ โโโโโโโโโโโโโโโโโโโโโโ MPerBlock
โ โ โ โ โโโโโโโโโโโโโโโโโโโโโโโโโโโโ BlockSize
โ โ โ โโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโ Scale_Block_K
โ โ โโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโ Scale_Block_N
โ โโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโโ Scale_Block_M (1=per-token)