Converts cuTile GPU kernels (@ct.kernel) to Triton (@triton.jit). Handles standard in-repo conversion, debugging (cudaErrorIllegalAddress, shape mismatch, numerical mismatch), and mapping cuTile idioms (ct.load/ct.store, ct.Constant, ct.launch) to Triton equivalents. Covers dual-kernel layout flags

cuTile → Triton Conversion Convert @ct.kernel kernels to @triton.jit . API mapping: [references/api mapping.md](./references/api mapping.md) (cuTile → Triton). In this skill’s Markdown, Triton launch syntax kernel［grid］(…) uses Unicode brackets so link checkers do not parse [grid](…) as a hyperlink; use normal ASCII brackets in real Triton code. Instructions Follow the phase gated workflow in [translations/workflow.md](./translations/workflow.md). Every conversion should go through analyze → convert → validate → test → benchmark , with explicit gates before moving on. Use the documents in [Workflow Selection]( workflow selection) when the task matches a special case (errors, layout flags, perf). 0. Optimization strategy (perf sensitive / attention) — If the op is attention, FMHA, sliding window, soft cap, or GQA (e.g. Gemma gemma attention ), read [references/optimization strategy.md](./references/optimization strategy.md) before converting the inner loop, then apply [§4 Gemma FMHA checklist](./references/optimization strategy.md 4 gemma fmha gemma attention conversion checklist mandatory) . For other GEMM/BMM/attention adjacent kernels, still skim §2–§3 of that file after TMA is done. 1. Select path — Existing TileGym op: standard mode in translations/workflow.md . If the cuTile source uses transpose / transpose v , dual layouts, or MLA style paths, read [translations/advanced patterns.md](./translations/advanced patterns.md) before writing Triton (two kernels + META grid, not one kernel + tl.trans ). 2. Pre flight — Run the [Pre flight Analysis]( pre flight analysis run before converting) grep commands on the cuTile source. Count @ct.kernel definitions; note TMA relevant ct.load / ct.store , ct.launch , Constant , and layout flags. 3. Read mapping — Keep [references/api mapping.md](./references/api mapping.md) open for cuTile → Triton API pairs. For runtime failures (illegal address, dtype, strides), use [references/debugging.md](./references/debugging.md). 4. Convert — Copy the [Conversion Checklist]( conversion checklist) into a todo list and execute in order. Structure and file placement: [translations/file structure.md](translations/file structure.md). Mandatory: any 2D+ block shaped tile load/store uses tl.make tensor descriptor (TMA), not raw tl.load(ptr+offs, mask=…) for full tiles—skipping this is the most common source of large regressions. Host side: Triton bracket launch <code kernel［grid］(args)</code with tuple or lambda META: (…) for autotune; no ct.launch . 5. Validate — Syntax check the new Triton module; run the relevant TileGym pytest targets for the op: pytest tests/ops/test <op .py k "triton" vs . Fix failures before benchmarking. 6. Benchmark — Compare Triton vs cuTile on perf tests. If Triton is clearly slower, follow PERFORMANCE ANALYSIS (Phase c2t 5) in [translations/workflow.md](./translations/workflow.md) and [references/optimizing reference.md](./references/optimizing reference.md) for GEMM/BMM/attention; use [references/optimization strategy.md](./references/optimization strategy.md) as the ordered checklist. If you see 10–50× slowdowns, read CRITICAL PERFORMANCE PATTERNS in that same workflow file first. Execution rules (MUST): Create and track the conversion checklist (e.g. TodoWrite) before editing kernel code; complete steps in order—do not skip pre flight or TMA decisions. For attention / FMHA / Gemma / GQA / soft cap / sliding window : read [references/optimization strategy.md](./references/optimization strategy.md) and apply §4 before treating the conversion as optimized. Do not ship raw pointer+mask 2D+ tile loads where TMA applies; document any intentional exception. If tests or benchmarks fail a gate, stop and fix before declaring the conversion done—do not stack unverified changes. Workflow Selection Existing TileGym op → Standard Mode: [translations/workflow.md](./translations/workflow.md) Errors ( cudaErrorIllegalAddress , shape mismatch, numerical mismatch) → [references/debugging.md](./references/debugging.md) Advanced patterns (TMA, dual layout flags transpose , autotune + META grid, Array.slice, ct.gather().item()) → [translations/advanced patterns.md](./translations/advanced patterns.md) (MLA style two kernels, avoid 3–15× regression on transpose=False ). Performance (Triton kernel slower than cuTile, autotuning, profiling) → [translations/workflow.md](./translations/workflow.md) (section PERFORMANCE ANALYSIS (Phase c2t 5) ) Optimization strategy hub (ordered checklist: advanced patterns + optimizing reference) → [references/optimization strategy.md](./references/optimization strategy.md) — read first for attention/FMHA/Gemma; then drill into the two source docs as needed Optimizing GEMM/BMM/attention (after TMA, or Triton 10–20% slower) → [references/optimizing reference.md](./references/optimizing reference.md) — EVEN K fast path, transpose via pointer arithmetic, grid layout, autotune breadth, epilogue subtile; use these patterns during conversion and before perf sign off (summarized in optimization strategy §2–§3 ) Gemma attention / GQA FMHA conversion → [references/optimization strategy.md §4](./references/optimization strategy.md 4 gemma fmha gemma attention conversion checklist mandatory) Blackwell optimization (complex kernels with iterative algorithms, register pressure, loop unrolling) → [references/optimizing reference.md](./references/optimizing reference.md) §9 — TMA descriptors, loop unroll factor , occupancy autotuning, TMEM friendly block sizes, slab allocator, dual path kernel design ⚠️ 10 50x REGRESSION (catastrophic slowdown after conversion) → [translations/workflow.md](./translations/workflow.md) — section CRITICAL PERFORMANCE PATTERNS (AVOID 10 50x REGRESSION) ⚠️ Good perf on transpose=True only, collapse on transpose=False (or opposite) → [translations/advanced patterns.md](./translations/advanced patterns.md) — §1 Dual layout flag; two @triton.jit kernels + grid = lambda META: (... META["BLOCK H"] ...) Pre flight Analysis (Run BEFORE converting) Conversion Checklist Copy this checklist and track progress: Gotchas (Most Common Translation Errors) { gotchas most common translation errors} Comprehensive table of patterns that frequently break or regress when porting @ct.kernel to @triton.jit — mma accumulator, type cast, grid, TMA usage, dtype handling, layout flags, batched matmul, etc. See: [references/gotchas.md](./references/gotchas.md) — read this BEFORE writing the Triton kernel. Performance Gotchas (10 50x Regression Risk) { performance gotchas 10 50x regression risk} ⚠️ These cause CATASTROPHIC slowdowns. Check BEFORE benchmarking. Patterns and their impact: TMA vs raw ptr+mask (5 20×), autotune vs fixed tile sizes (2 3×), broadcast to + tl.dot (10 50×), extract slice chains (2 5×), and more. See: [references/performance gotchas.md](./references/performance gotchas.md) — full regression risk table. Full details: [translations/workflow.md](./translations/workflow.md) — section CRITICAL PERFORMANCE PATTERNS (AVOID 10 50x REGRESSION) . Full API mapping: [references/api mapping.md](./references/api mapping.md). Triton math dtype (erf/erfc/exp/log/sqrt) and the "don't substitute erf with tanh" pattern: [references/debugging.md](./references/debugging.md) — section Triton Math Function Dtype Requirements (CRITICAL) . Optimization strategy (hub) File: [references/optimization strategy.md](./references/optimization strategy.md) Summarizes [translations/advanced patterns.md](./translations/advanced patterns.md) (layout flags, dual kernels, autotune+ META , batched launch, Blackwell pointers) and [references/optimizing reference.md](./references/optimizing reference.md) (post TMA micro opts, §9) into §1–§3 plus a mandatory §4 Gemma FMHA checklist . Rule: For attention / FMHA / Gemma style conversions, open optimization strategy in the same session as workflow — do not rely on TMA alone for perf sign off. Reference Documents { reference documents} Read from cuTile → Triton perspective. Core files live in this skill under . Category Document Content Strategy [optimization strategy.md](./references/optimization strategy.md) Ordered hub: advanced patterns + optimizing reference; §4 Gemma FMHA mandatory checklist Workflows [translations/workflow.md](translations/workflow.md) Standard c2t conversion (phases + checklist) [translations/file structure.md](translations/file structure.md) Where to place Triton files when converting from cuTile [translations/advanced patterns.md](./translations/advanced patterns.md) Dual layout flags (transpose), autotune + META grid, MLA style two kernels API [api mapping.md](./references/api mapping.md) cuTile → Triton mapping [optimizing reference.md](./references/optimizing reference.md) GEMM/BMM/attention optimizations (EVEN K, transpose, grid, autotune, epilogue subtile) Gotchas [gotchas.md](./references/gotchas.md) Common cuTile→Triton translation errors (mma, dtype, grid, TMA, layout flags) [performance gotchas.md](./references/performance gotchas.md) 10 50× regression risk table (TMA vs ptr+mask, broadcast to, extract slice chains, autotune) Testing & errors [references/debugging.md](./references/debugging.md) Triton runtime errors (cudaErrorIllegalAddress, pointer type, stride overflow) Worked Examples Use cutile kernel.py as source and triton kernel.py as target : Example Directory Complexity Vector Add [examples/01 vector add/](examples/01 vector add/) Basic Softmax [examples/02 softmax/](examples/02 softmax/) Intermediate LayerNorm [examples/03 layernorm/](examples/03 layernorm/) Intermediate MatMul [examples/04 matmul/](examples/04 matmul/) Advanced Attention [examples/05 attention/](examples/05 attention/) Advanced Read cutile kernel.py first, then triton kernel.py , to see the inverse mapping. ⚠️ MANDATORY COMPLETION CHECKLIST (DO NOT SKIP) A conversion is NOT COMPLETE until ALL items are checked. Copy and complete: Why this matters: Gate 1 catches functional bugs Gate 2 prevents catastrophic 5 20x regressions (most common mistake) Gate 3 validates that optimization was effective Gate 4 creates accountability record If any gate fails: Fix and re verify before declaring complete.

tilegym-converting-cutile-to-triton