/tilegym-converting-cutile-to-triton

cuTile to Triton Conversion Workflow

Guide for converting cuTile kernels to Triton with the same rigor as the inverse (Triton→cuTile).

🚀 TODO WORKFLOW (MANDATORY - CREATE IMMEDIATELY)

Upon starting a cuTile→Triton conversion task, IMMEDIATELY create this todo list using todowrite:

todowrite([
  { id: "c2t-1", content: "[Optional] Analyze test coverage - verify cuTile tests pass, identify edge cases", status: "pending", priority: "medium" },
  { id: "c2t-2", content: "Convert cuTile → Triton - apply API mapping, generate Triton file", status: "pending", priority: "high" },
  { id: "c2t-3", content: "Test correctness - run python -m pytest -k 'triton', fix errors (max 5 attempts)", status: "pending", priority: "high" },
  { id: "c2t-4", content: "TMA optimization (MANDATORY) - replace ALL 2D+ raw ptr+mask with tl.make_tensor_descriptor; add tl.assume() alignment hints; add autotuning; if transpose/layout flag, two kernels + lambda META grid (advanced-patterns.md); skip = 5-20x regression", status: "pending", priority: "high" },
  { id: "c2t-5", content: "Performance test - run pytest -k 'test_perf' --print-record, compare vs cuTile, optimize if >20% slower", status: "pending", priority: "high" }
])

Workflow Execution Rules

Phase → Todo Mapping

DO NOT ask "should I proceed?" - execute the full workflow end-to-end.

RATIONALE: Key Thresholds

Why these values? (Aligned with Triton→cuTile skill.)

Relaxed tolerances: Use 2× for reductions, transcendentals, chained ops.

VALIDATION LOOP (MANDATORY)

NEVER proceed until tests pass. This pattern applies to ALL test phases.

┌─────────────────────────────────────────────────────────┐
│                   VALIDATION LOOP                        │
│                                                          │
│   ┌─────────┐     ┌─────────┐     ┌─────────┐          │
│   │  RUN    │────▶│  CHECK  │────▶│  PASS?  │          │
│   │  TEST   │     │  OUTPUT │     │         │          │
│   └─────────┘     └─────────┘     └────┬────┘          │
│        ▲                               │               │
│        │              ┌────────────────┼───────────┐   │
│        │              │                │           │   │
│        │              ▼                ▼           │   │
│   ┌─────────┐    ┌─────────┐     ┌─────────┐      │   │
│   │  FIX    │◀───│   NO    │     │  YES    │──────┘   │
│   │  ERROR  │    │(attempt │     │  DONE   │          │
│   └─────────┘    │  < 5)   │     └─────────┘          │
│                  └─────────┘                          │
└─────────────────────────────────────────────────────────┘

Validation Checklist (copy for each attempt):

- [ ] Attempt #__: Run test command
- [ ] Check: `X passed, 0 failed`?
- [ ] No `FAILED tests/...` markers?
- [ ] No exceptions (syntax, shape, dtype)?
- [ ] If FAIL: identify error → fix → increment attempt
- [ ] If attempt >= 5: STOP, escalate to user

CRITICAL: Do NOT proceed to next phase until loop completes successfully.

EXACT TEST COMMANDS (LOW FREEDOM)

DO NOT MODIFY these commands. Flags are validated for correct output handling.

Correctness Test (Triton)

# DO NOT MODIFY
python -m pytest {test_path} -k "test_op and triton" -vs --tb=short 2>&1 | tail -100

Performance Test

# DO NOT MODIFY
python -m pytest {test_path} -k "test_perf" --print-record -v 2>&1 | tail -50

Triton Debug / Profiling

TRITON_INTERPRET=1 python script.py
TRITON_PRINT_AUTOTUNING=1 python script.py
TILEIR_DUMP_DIR=/tmp/dumping/triton python -m pytest {test_path} -k "test_op and triton" --timeout=120

Flag rationale: -vs (verbose + no capture), --tb=short (concise tracebacks), 2>&1 | tail -100 (capture stderr, limit for context).

TABLE OF CONTENTS

🚀 TODO WORKFLOW (MANDATORY - CREATE IMMEDIATELY) RATIONALE: Key Thresholds VALIDATION LOOP (MANDATORY) EXACT TEST COMMANDS (LOW FREEDOM) MODE SELECTION COMMAND CHEAT SHEET TMA OPTIMIZATION (Phase c2t-4) PERFORMANCE ANALYSIS (Phase c2t-5) CRITICAL PERFORMANCE PATTERNS DUAL LAYOUT FLAG + AUTOTUNE GRID (MLA-style) MEMORY LAYOUT PATTERNS Standard Conversion Steps Quick Reference: cuTile → Triton

MODE SELECTION

Standard Mode (cuTile → Triton)

Use when: Converting existing TileGym cuTile operators to Triton (e.g. for portability or comparison).

Path convention: /cutile/ → /triton/

TRITON_PATH="${CUTILE_PATH//\/cutile\//\/triton\/}"
mkdir -p $(dirname $TRITON_PATH)

COMMAND CHEAT SHEET

# Standard mode path derivation
TRITON_PATH="${CUTILE_PATH//\/cutile\//\/triton\/}"
mkdir -p $(dirname $TRITON_PATH)

# Correctness testing
python -m pytest {test_path} -k "test_op and triton" -vs

# Performance testing (Triton vs cuTile)
python -m pytest {test_path} -k "test_perf and (triton or cutile)" --print-record -v

# Triton profiling / autotune visibility
TRITON_PRINT_AUTOTUNING=1 python -m pytest {test_path} -k "test_op and triton"
triton.testing.do_bench(lambda: kernel［grid］(launch_args))  # In script: measure ms

TMA OPTIMIZATION (Phase c2t-4) {#tma-optimization-phase-c2t-4}

This phase is MANDATORY. Do not skip. Raw pointer + mask for 2D+ tile loads are 5-20x (500%-2000%) slower than TMA on Hopper/Blackwell. Converted kernels that use only tl.load(ptr+offs, mask=...) for block-shaped 2D+ access will show severe regression until TMA is added.

When TMA is Required

• 2D+ tile loads (GEMM, attention, convolution, any load with shape like (BLOCK_M, BLOCK_K)) → ALWAYS use TMA (tl.make_tensor_descriptor + .load([...]))
• 1D loads (single contiguous block, elementwise, simple reductions) → raw pointer OK

TMA Conversion Pattern

# BEFORE (cuTile ct.load) - already uses TMA internally
tile = ct.load(arr, index=(bid_m, bid_k), shape=(BLOCK_M, BLOCK_K))

# WRONG (naive Triton conversion) - 10-20x SLOWER
offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
offs_k = tl.arange(0, BLOCK_K)
tile = tl.load(ptr + offs_m[:, None] * stride_m + offs_k[None, :], mask=mask)

# CORRECT (TMA tensor descriptor) - FAST
desc = tl.make_tensor_descriptor(
    base=ptr,
    shape=[M, K],
    strides=[stride_m, 1],
    block_shape=[BLOCK_M, BLOCK_K],
)
tile = desc.load([pid_m * BLOCK_M, pid_k * BLOCK_K])

TMA Setup (Required Once)

from typing import Optional

def alloc_fn(size: int, alignment: int, stream: Optional[int]):
    return torch.empty(size, device="cuda", dtype=torch.int8)

triton.set_allocator(alloc_fn)

TMA Checklist (Run BEFORE moving to c2t-5) {#tma-checklist-run-before-moving-to-c2t-5}

TMA Optimization Verification:
 [ ] All 2D+ tile loads use tl.make_tensor_descriptor
 [ ] Masks removed (TMA handles bounds automatically)
 [ ] tl.assume() alignment hints added for strides: tl.assume(stride % 8 == 0)
 [ ] tl.assume() alignment hints added for pointers: tl.assume(ptr.to(tl.int64) % 16 == 0)
 [ ] Autotuning added with GPU-specific configs (see PERFORMANCE ANALYSIS)

PERFORMANCE ANALYSIS (Phase c2t-5) {#performance-analysis-phase-c2t-5}

When Triton is >20% slower than cuTile, follow this systematic flow.

Step 1: Benchmark

python -m pytest {test_path} -k "test_perf and (triton or cutile)" --print-record -v

If Triton is within 20% of cuTile → done. If not → continue.

Step 2: Match Tile Sizes and Grid

• Constexpr block sizes — Use the same BLOCK_M, BLOCK_N, BLOCK_K (or BLOCK_SIZE) as the cuTile kernel as tl.constexpr in Triton.
• Grid — Use triton.cdiv(M, BLOCK_M) etc. so grid shape matches cuTile's (n_m, n_n, 1)-style launch.

Step 3: Memory Access Pattern

• cuTile uses block index in ct.load(arr, index=(...), shape=(...)).
• Triton uses element offset: offs = pid * BLOCK + tl.arange(0, BLOCK) (and strides for 2D+). Ensure tl.load(ptr + offs, mask=...) or tl.make_block_ptr + load matches the same coalescing and alignment.
• Cache hints — For memory-bound kernels, try tl.load(..., cache_modifier=".cg") if appropriate (see performance-model.md).

Step 4: Triton Autotuning

If the cuTile op uses autotune (tile sizes, occupancy), add triton.autotune to the Triton kernel so Triton can search the same space:

@triton.autotune(
    configs=[
        triton.Config({'BLOCK_M': 128, 'BLOCK_N': 256, 'BLOCK_K': 64}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_M': 64, 'BLOCK_N': 256, 'BLOCK_K': 32}, num_stages=4, num_warps=4),
        # ... match or expand cuTile config space
    ],
    key=['M', 'N', 'K'],  # Retune when these change
)
@triton.jit
def kernel(a_ptr, b_ptr, c_ptr, M, N, K, ...,
           BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    ...

Triton autotune knobs: BLOCK_* (tile dimensions), num_warps (1, 2, 4, 8, 16, 32), num_stages (2–5 typical).

Step 5: Profiling and Bottleneck

• Run with TRITON_PRINT_AUTOTUNING=1 to see which config is chosen.
• Use triton.testing.do_bench(lambda: kernel［grid］(launch_args)) for stable timings.
• For deeper analysis: Nsight Compute, or dump Triton IR (TILEIR_DUMP_DIR=/tmp/dumping/triton) and compare with cuTile IR (see references/debugging.md) if the gap remains unexplained.

Step 6: tf32 and Dtypes

• Triton defaults allow_tf32=True for tl.dot. If cuTile used an explicit tf32 cast before ct.mma, behavior should already match; use allow_tf32=False only if you need strict fp32.

Bottleneck checklist: algorithmic/grid → memory access → occupancy/warps → micro-optimizations.

Standard Conversion Steps

Quick Reference: cuTile → Triton

Full cuTile → Triton API mapping: references/api-mapping.md.

CRITICAL PERFORMANCE PATTERNS (AVOID 10-50x REGRESSION) {#critical-performance-patterns-avoid-10-50x-regression}

Case study: Group GEMM conversion showed 20-50x regression. Root causes and fixes below.

Performance Killer #1: Raw Pointer Arithmetic vs TMA Tensor Descriptors

Impact: 5-20x slowdown (500%-2000% regression) — most common cause of conversion regression

TMA (Tensor Memory Accelerator) enables async bulk memory transfers on Hopper/Blackwell. Raw pointer arithmetic with masks for 2D+ tile loads is dramatically slower. Always use TMA for 2D+ block-shaped loads; do not ship conversion with raw ptr+mask for GEMM/attention/conv tiles.

# SLOW (10-20x regression) - Raw pointer + masks
offs_m = tile_m * TILE_M + tl.arange(0, TILE_M)
offs_k = tl.arange(0, TILE_K)
a_ptrs = A_ptr + (offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak)
a_mask = (offs_m[:, None] < M) & (offs_k[None, :] < K)
a = tl.load(a_ptrs, mask=a_mask, other=0.0)  # Masked load = SLOW

# FAST - TMA tensor descriptors
a_desc = tl.make_tensor_descriptor(
    base=a_base_ptr,
    shape=[m, k],
    strides=[lda, 1],
    block_shape=[BLOCK_M, BLOCK_K],
)
a = a_desc.load([offset_am, kk * BLOCK_K])  # Bulk TMA load = FAST

When to use TMA:

• Blackwell (sm_120) or Hopper (sm_90) GPUs
• Matrix operations (GEMM, attention, convolution)
• Any kernel with 2D+ tile loads

TMA setup requirements:

from typing import Optional

# TMA allocator (required once per kernel launch context)
def alloc_fn(size: int, alignment: int, stream: Optional[int]):
    return torch.empty(size, device="cuda", dtype=torch.int8)

triton.set_allocator(alloc_fn)

Performance Killer #2: Inefficient Group/Batch Iteration

Impact: 2-5x slowdown

For grouped operations (group GEMM, batched attention), how you find which group a tile belongs to matters.

# SLOW (2-5x regression) - Linear search ALL groups per tile
@triton.jit
def _find_group_id(tile_idx, problem_offsets, num_groups: tl.constexpr):
    group_id = num_groups - 1
    for g in range(num_groups):  # Scans ALL groups every time
        offset_start = tl.load(problem_offsets + g)
        offset_end = tl.load(problem_offsets + g + 1)
        is_in_group = (tile_idx >= offset_start) & (tile_idx < offset_end)
        group_id = tl.where(is_in_group, group_id, g)  # Conditional per group
    return group_id

# FAST - Natural while-loop advancement
@triton.jit
def kernel(...):
    tile_idx = tl.program_id(0)
    last_problem_end = 0

    for g in range(group_size):
        # ... load group dimensions ...
        num_tiles = num_m_tiles * num_n_tiles

        # Only process tiles belonging to this group
        while tile_idx >= last_problem_end and tile_idx < last_problem_end + num_tiles:
            # Process tile
            tile_idx += num_programs  # Persistent scheduling

        last_problem_end += num_tiles  # Advance boundary

Performance Killer #3: Missing Autotuning

Impact: 2-3x slowdown

Fixed tile sizes vs architecture-optimized configurations.

# SLOW - Fixed sizes, no tuning
TILE_M, TILE_N, TILE_K = 128, 128, 64  # May be wrong for your GPU/problem

# FAST - Autotuning with GPU-specific configs
def _get_configs():
    gpu_cap = torch.cuda.get_device_capability()
    if gpu_cap in [(12, 0), (12, 1)]:  # Blackwell
        return [
            triton.Config({"BLOCK_M": BM, "BLOCK_N": BN, "BLOCK_K": BK})
            for BM in [64, 128]
            for BN in [64, 128, 256]
            for BK in [64, 128]
        ]
    elif gpu_cap == (9, 0):  # Hopper
        return [
            triton.Config({"BLOCK_M": BM, "BLOCK_N": BN, "BLOCK_K": BK},
                         num_stages=s, num_warps=w)
            for BM in [128, 256]
            for BN in [128, 256]
            for BK in [64, 128]
            for s in [4, 5]
            for w in [8]
        ]
    # ... other architectures

@triton.autotune(configs=_get_configs(), key=["group_size", "dtype"])
@triton.jit
def kernel(...):
    ...

Performance Killer #4: Missing Alignment Hints

Impact: 1.5-2x slowdown

Triton compiler can optimize better with alignment guarantees.

# SLOW - No hints, compiler assumes worst case
lda = tl.load(group_strides + g * 3)
a_base_ptr = tl.load(group_a_ptrs + g).to(tl.pointer_type(dtype))

# FAST - Alignment hints enable compiler optimizations
lda = tl.load(group_strides + g * 3)
tl.assume(lda % 8 == 0)  # Stride is 8-element aligned

a_base_ptr = tl.load(group_a_ptrs + g).to(tl.pointer_type(dtype))
tl.assume(a_base_ptr.to(tl.int64) % 16 == 0)  # 16-byte aligned pointer

Performance Killer #5: Unnecessary Masking

Impact: 1.2-1.5x slowdown

Masks add predication overhead. TMA handles boundaries automatically.

# SLOW - Masks on every operation
a = tl.load(a_ptrs, mask=a_mask, other=0.0)
b = tl.load(b_ptrs, mask=b_mask, other=0.0)
tl.store(c_ptrs, c, mask=c_mask)

# FAST - TMA handles bounds, no masks needed
a = a_desc.load([offset_am, kk * BLOCK_K])
b = b_desc.load([kk * BLOCK_K, offset_bn])
c_desc.store([offset_cm, offset_cn], c)

Performance Killer #6: K-loop Offset Recalculation

Impact: 1.1-1.2x slowdown

For GEMM K-loops, avoid recalculating full offsets each iteration.

# SLOW - Recalculate every iteration
for k_tile in range(num_k_tiles):
    k_offset = k_tile * TILE_K
    a_ptrs = A_ptr + (offs_m[:, None] * stride_am + (k_offset + offs_k[None, :]) * stride_ak)
    # ... full offset calculation each time

# FAST - Increment pointers (for non-TMA kernels)
a_ptrs = a_ptr + offs_am[:, None] * lda + offs_k[None, :]
for kk in range(0, tl.cdiv(k, BLOCK_K)):
    tl.multiple_of(a_ptrs, [16, 16])  # Pipeline hint
    a = tl.load(a_ptrs)
    a_ptrs += BLOCK_K  # Simple increment

Performance Checklist (Run BEFORE declaring conversion complete)

Performance Verification:
 [ ] TMA tensor descriptors used for 2D+ tile loads (Hopper/Blackwell)
 [ ] Autotuning added with GPU-specific configs
 [ ] Group/batch iteration uses natural loop advancement (not linear search)
 [ ] tl.assume() hints added for stride and pointer alignment
 [ ] Masks removed where TMA handles boundaries
 [ ] K-loop uses pointer increment or TMA offset (not full recalc)
 [ ] Pipeline hints (tl.multiple_of) added for non-TMA loads
 [ ] num_stages and num_warps tuned in autotune configs
 [ ] Benchmark shows <20% regression vs original cuTile
 [ ] Memory layout matches original (transpose + contiguous, NOT 5D reshape)
 [ ] Dimension params marked as tl.constexpr (bs, hd, n_heads, pad_*)
 [ ] No unnecessary tensor clones (transpose + contiguous suffices)
 [ ] If `transpose`/layout flag: two kernels + `grid = lambda META: (triton.cdiv(..., META["BLOCK_H"]), ...)` — see [advanced-patterns.md](advanced-patterns.md)

DUAL LAYOUT FLAG + AUTOTUNE GRID (MLA-style)

When: cuTile or the public API exposes transpose (or equivalent) and perf tests show one mode fast, the other 3–15× slow.

Do not implement the non-transpose case by reusing the transpose kernel with extra tl.trans on qk / p and a [BLOCK_N, BLOCK_H] softmax state unless that matches a proven-fast reference.

Do implement two separate kernels following this structure:

Autotune: If @triton.autotune supplies BLOCK_H / BLOCK_N, the host must not pass fixed blocks into torch.autograd.Function.apply. Use:

grid = lambda META: (triton.cdiv(num_head, META["BLOCK_H"]), B, 1)
kernel with grid then call (tensor_bases…, BLOCK_D=BLOCK_D, BLOCK_KPE=BLOCK_KPE)

Full pattern, diagnosis table, and host-descriptor caveat: advanced-patterns.md.

MEMORY LAYOUT PATTERNS (AVOID 50-150% REGRESSION)

Case study: RoPE conversion showed 50-150% regression. Root cause: wrong memory layout.

Performance Killer #7: 5D Reshape vs Contiguous Access

Impact: 50-150% slowdown

When kernels access tensor halves (first [0:dim//2], second [dim//2:dim]), naive 5D reshape breaks memory coalescing.

# SLOW (50-150% regression) - 5D reshape
# Original cuTile might reshape: q.reshape(bsz, n_head, seq_len, 2, head_dim//2)
# Then access via stride_2 dimension - NON-CONTIGUOUS
q_offs_1 = batch * q_stride_b + heads[:, None] * q_stride_h + 0 * q_stride_2 + hd[None, :]
q_offs_2 = batch * q_stride_b + heads[:, None] * q_stride_h + 1 * q_stride_2 + hd[None, :]

# FAST - Transpose + contiguous + offset arithmetic
# Transpose to: (bsz, seq_len, n_head, head_dim) - head_dim CONTIGUOUS
q = q.transpose(1, 2).contiguous()
q = q + pid * q_row_stride  # Row-based addressing
first_half_offs = heads[:, None] * hd + tl.arange(0, pad_hd // 2)[None, :]
second_half_offs = first_half_offs + (hd // 2)  # Just add offset!
q_tile_1 = tl.load(q + first_half_offs, mask=mask)
q_tile_2 = tl.load(q + second_half_offs, mask=mask)

Why this matters:

• 5D reshape with stride_2 creates non-sequential access: [..., 0, :half_hd] then [..., 1, :half_hd]
• Transpose + offset gives sequential access: [:half_hd] then [half_hd:] - both contiguous
• Memory coalescing difference: ~50-150% performance gap

Performance Killer #8: Missing constexpr on Dimension Parameters

Impact: 10-20% slowdown

Triton compiler can optimize better when dimensions are compile-time constants.

# SLOW - All parameters dynamic
def kernel(q, k, n_qh, n_kh, hd, pad_hd, BACKWARD_PASS):
    # Compiler cannot unroll loops or optimize register allocation

# FAST - Dimension params as constexpr
def kernel(q, k,
           n_qh: tl.constexpr,
           n_kh: tl.constexpr,
           hd: tl.constexpr,
           pad_hd: tl.constexpr,
           bs: tl.constexpr,
           BACKWARD_PASS: tl.constexpr = False):
    # Compiler can unroll, specialize, and optimize

Performance Killer #9: Unnecessary Tensor Clones

Impact: 10-20% slowdown

In-place operations don't need explicit clones when transpose + contiguous already creates a copy.

# SLOW - Explicit clone overhead
def forward(ctx, q, k, ...):
    q = q.clone()  # Unnecessary copy
    k = k.clone()  # Unnecessary copy
    q, k = rope_forward(q, k, ...)

# FAST - Transpose + contiguous is sufficient
def rope_forward(q, k, ...):
    q = q.transpose(1, 2)  # Returns view
    k = k.transpose(1, 2)  # Returns view
    q = q.contiguous()     # Creates copy (if needed)
    k = k.contiguous()     # Creates copy (if needed)
    # Now safe to modify in-place

Pre-flight: Detect Split-Dimension Patterns

# Check if kernel splits a dimension
grep "reshape.*2," source.py        # Reshape with 2 in shape
grep "stride_2\|stride(3)" source.py  # Extra stride dimension
grep "half_\|// 2" source.py        # Half-dimension patterns

If found, use transpose + offset pattern, NOT 5D reshape.

Quick Diagnosis: Why Is My Converted Kernel Slow?