Convert a GPU-shaped Triton kernel into a full Spyre-aware kernel (spyre.py) — tensor descriptors plus the authoring invariants (see _shared/invariants/), scratchpad batching, and Spyre-compiler gap handling. Use when asked to make a kernel Spyre-aware or port to Spyre. For a plain raw->descriptor port with no invariants, use td-convert.
Convert a GPU-shaped Triton kernel into a full Spyre-aware kernel (spyre.py) — tensor descriptors plus the authoring invariants (see _shared/invariants/), scratchpad batching, and Spyre-compiler gap handling. Use when asked to make a kernel Spyre-aware or port to Spyre. For a plain raw->descriptor port with no invariants, use td-convert.
Spyre Kernel Conversion Skill
Convert a GPU-shaped Triton kernel (original.py) into a Spyre-aware kernel
(spyre.py) that satisfies the authoring invariants for IBM Spyre/AIU.
This builds on the tensor-descriptor conversion (td-convert): first the
memory accesses become descriptors, then this skill adds the Spyre-specific
distribution, scratchpad, and gap handling. If you only need the descriptor port
(no invariants), use td-convert.
Trigger
Use when asked to convert a kernel to Spyre-aware form, port a kernel to Spyre,
or produce a spyre.py from an original.py.
Every Spyre-aware kernel must satisfy the invariants in
../_shared/invariants/ — one file each. Read that
directory and apply every file you find; the set may grow or shrink, so do not
rely on a fixed list here.
Procedure
Descriptors first. Apply
../_shared/descriptor-rules.md in full —
replace pointer arithmetic, drop redundant masks, respect the 16-byte
minimum. (Same as td-convert step 1–3.)
Satisfy every invariant. Read
../_shared/invariants/ and apply each file's rule
and canonical pattern — do not work from a fixed list, as the set may change.
Typical work this entails: replacing the unbounded GPU grid with a
tl.program_id / tl.num_programs distribution loop, using fixed constexpr
tiles with a loop over the dim (never BLOCK = next_power_of_2(N)), and
keeping problem sizes as runtime args with tl.cdiv / tl.minimum tail
handling. The invariant files are authoritative for the details.
Strip @triton.autotune and other GPU-only helpers (tl.assume,
tl.multiple_of, CUDA/HIP config). Why autotune specifically: on GPU,
autotune times each BLOCK_* config and caches the fastest per key — it
relies on GPU-style tight kernel-launch timing. Spyre does not work that way:
tile sizes interact with the 2 MB scratchpad budget and OpSpec-level
scheduling (LoopSpec, tiled_symbols, device_size), so the right sizes come
from compiler-level analysis, not random search. (Read this as a hard
constraint, not a style preference.)
Preserve the signature when possible so the wrapper can call both forms
with minimal branching. Recover grid-derived values from existing args
before adding parameters; update the wrapper if you must add any.
Utilize the scratchpad — batch work items. After the single-item
conversion, evaluate whether the distribution loop processes one item per
iteration with a small accumulator. If so, batch BLOCK_ITEMS per iteration
per ../_shared/spyre/scratchpad-batching.md.
This also documents the 16-byte escape (a scalar load widened to a
[BLOCK_ITEMS] vector clears the minimum) and the divergent-control-flow
masking trade-off.
Handle Spyre-compiler descriptor gaps the right way. Write the cleanest
descriptor-first form even if it does not compile yet; if it hits a gap,
mark the site with a # [gap] annotation rather than baking a workaround
into the kernel. See ../_shared/spyre/gap-handling.md
for the annotation contract and the live sources (KB, torch-spyre/triton
issues/docs) where the current gaps are tracked, so the target form is clear
once the compiler catches up.
Layout awareness (write logical; the compiler tiles physical)
Write descriptors in logical shape — the tensor's math dimensions — see the
worked ../_shared/examples/matmul-logical.md.
You do not hand-write the physical device layout; the compiler derives it.
Output file structure
# SPDX-License-Identifier: <same as original># SPDX-FileCopyrightText: <same as original>## Spyre-aware conversion of <kernel_function_name>.# Original: <source path># Changes summarized in kernels/<name>/conversion-notes.md.import triton
import triton.language as tl
@triton.jitdef_<name>_kernel_spyre(..., BLOCK_...: tl.constexpr):
"""<Brief description.>"""
pid = tl.program_id(0)
num_cores = tl.num_programs(0)
# descriptors ...# distribution loop ...
Document the conversion
After writing the kernel, record what changed in
kernels/<name>/conversion-notes.md — the single file shared with td-convert.
Append a ## Spyre-aware conversion section; if td-convert already wrote a
## Tensor-descriptor conversion section, leave it intact and add yours below.
Create the file with both headings if it does not exist yet.
# <name> conversion notes## Tensor-descriptor conversion- (written by td-convert, if it ran — leave intact)
## Spyre-aware conversion- Source: original.py → spyre.py
- Invariants: <keyper-invariantdecisions, e.g.whichaxisisthedistributionloop>- Scratchpad batching: <BLOCK_ITEMSchoiceandwhy, or "none — alreadydense">- Gaps: <gapsites + `# [gap]` annotation, or "none">-<anysignaturechangeorothernon-obviousdecision>
Keep it short — the why behind non-obvious choices, not a line-by-line diff.
Checklist
All memory accesses use descriptors (no raw tl.load/tl.store,
make_block_ptr, advance)
Scalar descriptors resolved (last dim ≥ 16 bytes) or kernel declared
non-portable, per descriptor-rules.md §4 — not annotated and shipped.
Spyre-compiler gaps written in clean form + # [gap] annotation — not
worked around (see gap-handling.md)
@triton.autotune removed, with the why understood (compiler-level
tiling, not launch-timing search)
Every invariant in _shared/invariants/ satisfied (read the directory —
do not assume a fixed set)
Scratchpad utilization evaluated; batched where beneficial
Output at kernels/<name>/spyre.py, kernel _<name>_kernel_spyre
kernels/<name>/conversion-notes.md updated with the Spyre-aware section
