Write or review Triton kernels for vLLM, with practical guidance for generated-code inspection, launch grids, indexing, specialization, tuning, and representative performance validation.
SkillScore82Maintained
Automatic
◆AI Engineer pack
Last verified 10/7/2026
What it changes
Passed 18 of 22 test cases with the skill, and 11 of 22 without it.
Without skill50%
With skill82%
Checks it passed that the model alone missed
The response specifies placing alternating integer scalar flags in do_not_specialize.
The response refers specifically to the official Triton semantics documentation page.
The response recommends using $kernel-microbenchmark for benchmark measurement and reporting.
and 6 more in the full report
7 fixed11 passed either way4 still failing0 broke
Pass-rate lift
+32 pp
Tokens
+14%
Agent turns
0%
Response time
-12%
Model gemini-3.6-flashMeasured 10/7/2026v118 of 22 cases passed with skill
Tokens, turns and time describe resource use, not a dollar price.
See it in action
One of 7 test cases the skill fixed. Same model, same prompt, with and without the skill.
1 of 7
Prompt
Our Triton kernel produces unexpected results when calculating negative array indices using Python-style modulo operations like (index - pad) % stride. Where in the official Triton documentation should we check for differences between Triton and Python/NumPy semantics?
What was checked, without and with the skill
Checked
Without skill
With skill
The response refers specifically to the official Triton semantics documentation page.
Missed
Passed
The response notes that integer modulo and division behavior in Triton can differ from Python defaults.
Passed
Passed
case-20 · gemini-3.6-flash · 10/7/2026 · answers are capped and redacted
What's inside
The instructions your model reads when the skill loads.
SKILL.md57 linessingle file
Triton Kernel Writing
Implementation
Follow the official
Triton semantics.
Check it when behavior may differ from Python or NumPy, especially type
promotion, integer division and modulo, casts, broadcasting, and variable
scoping.
Use the Triton kernel generated by torch.compile as a possible
implementation to inspect. Print Inductor's generated code with
TORCH_LOGS="output_code" .venv/bin/python <script> or enable
torch._logging.set_logs(output_code=True) before the compiled function
runs. Treat generated code as a reference, not as proof of correctness or
optimality.
Find reasonable defaults for compile-time knobs such as BLOCK_SIZE, or use
a small, legible heuristic when workloads need different choices. Use
triton.autotune only when tuning is critical to performance, such as for a
matrix multiplication. Otherwise prioritize simple code and fast startup.
Be careful to avoid unintended runtime JIT compilation. For example, put
unimportant runtime integer scalars in do_not_specialize, especially those
that may alternate between values such as 0 and 1, which can produce
different specialization keys.
The Triton compiler does not guarantee safe ordering when a kernel writes to
a pointer and subsequently reads from the same pointer. This pattern must
have a tl.debug_barrier() between the write and read. The barrier
synchronizes threads in the block; it does not synchronize separate program
instances.
Launch and Indexing
grid[1] and grid[2] must be at most 65,535. Choose or flatten the grid
order so those dimensions cannot exceed the limit for supported shapes.
For example, num_tokens is commonly 8K or 16K, but users may configure 32K
or more. If num_tokens is a grid dimension, it is safe to put it in
grid[0] (or tile it).
Use int64 for offset arithmetic when an index can exceed 32-bit range,
especially for KV-cache addressing. Cast operands before multiplication or
addition so an intermediate does not overflow in 32-bit arithmetic.
A [num_tokens, num_heads] grid can be a good low-latency mapping for decode,