Capstone kernel
The capstone kernel must combine correctness, a measured speedup, and a bottleneck story that survives profiling.
Three artifacts, one kernel
The capstone is a kernel plus three claims. Correctness: every benchmarked shape and dtype compared against a reference, with the tolerance justified by accumulation order and input dtype. Speed: a named baseline, named shapes, and <ms> and <speedup> with the band from one interleaved session. Bottleneck: a roofline prediction and the counter that confirms it.
Predict the counter before you profile
Write the prediction down before ncu runs. A kernel that is bandwidth-limited at <GB/s> should show dram__throughput.avg.pct_of_peak_sustained_elapsed high and sm__throughput.avg.pct_of_peak_sustained_elapsed low, and dram__bytes.sum should match the traffic the shape implies. If the counters disagree, the model was wrong, and the honest move is to fix the model and then the kernel. A session that only confirms what you already believed is decoration.
What survives
A negative result is a result. If the fused kernel loses at short sequence lengths because launch overhead dominates, say so and show the gap in the nsys timeline, where the kernel time is a sliver between gaps. If the win disappears against a baseline that is already fused, that is the finding, and it is worth more in an interview than a ratio against a straw man.
Things that break
Choosing the one shape where the kernel wins and presenting it as the kernel’s performance. A correctness check that runs on a single square case while the benchmark runs on six shapes. A bottleneck story contradicted by the counters that stays in the writeup because the roofline predicted it, which is the failure this node exists to prevent. And a speedup measured while tuning runs inside the timed region, so the number includes compilation. The harness uses Triton 3.9 do_bench, which returns milliseconds; confirm the return units against your installed version before the number enters the table.
Sources
- Nsight Compute Profiling Guidedocs
- Roofline: An Insightful Visual Performance Model for Multicore Architecturespaper
- triton.testing: do_bench and benchmark helpersdocs · 3.9
A version records the documentation the claim was checked against, not a house pin. Confirm flags with --help on the release you installed.
Check yourself
Your roofline says bandwidth-bound, but ncu reports DRAM at 20 percent and the tensor pipe near saturated. What now?
Which correctness check is adequate for the capstone?
The kernel is faster only at the largest shape. What does that support?