Menu
CUDA graph replay benchmark + source trace templates
Instalar com Codex ou Claude Copie este prompt, cole no Codex, Claude ou outro assistente e deixe que ele revise a pรกgina da skill e instale para vocรช.
Baseado na classificaรงรฃo ocupacional SOC
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)