/tilegym-converting-cutile-to-triton

cuTile → Triton API Mapping

Contents

Import & Decorator Indexing Memory Operations Tensor Creation Reductions Matrix Operations Type Operations Math Operations Comparison & Logic Bitwise Operations Debug Operations Atomic Operations Synchronization Host Functions Data Types Launch Patterns cuTile → Triton Gotchas Quick Reference Card TensorDescriptor Pattern (ct.load/ct.store → Triton TMA) Multi-dimensional Indexing Array.slice → Triton (Ragged Tensors) ct.gather().item() → Triton (Runtime Index TMA)

This document provides cuTile → Triton mappings for converting @ct.kernel code to @triton.jit. Source column is cuTile; target column is Triton.

Import & Decorator

Indexing

Memory Operations

ct.load / ct.store → Triton

⚠️ TMA for 2D+ loads: cuTile uses TMA internally for block-aligned 2D+ loads. Converting to raw tl.load(ptr + offs, mask=m) for 2D+ tile shapes causes 500%-2000% (5-20x) regression. For any load with 2D+ block shape (e.g. GEMM tiles, attention tiles), use TMA: tl.make_tensor_descriptor(...).load([...]) — see TensorDescriptor Pattern. Use raw ptr+mask only for 1D or truly scattered access.

cuTile uses block index in ct.load(arr, index=(...), shape=(...)). Triton uses element offset (ptr + offs), block ptr, or TMA tensor descriptor. When converting:

Gather/Scatter → Pointer load/store

Critical: In Triton, tl.load(ptr + offs, ...) uses element offset offs, not block index. Build offs from tl.program_id(axis) and tl.arange(0, BLOCK) (and strides).

Tensor Creation

Reductions

Matrix Operations

float32 → tf32 in Triton

In cuTile you may have:

a_mma = ct.astype(a, ct.tfloat32) if a.dtype == ct.float32 else a
b_mma = ct.astype(b, ct.tfloat32) if b.dtype == ct.float32 else b
acc = ct.mma(a_mma, b_mma, acc=acc)

In Triton, default behavior already matches:

# Triton: allow_tf32=True by default
acc = tl.dot(a, b, acc=acc)

Use allow_tf32=False only if you need strict IEEE float32.

Type Operations

Math Operations

Index Arithmetic

Comparison & Logic

Bitwise Operations

Debug Operations

# cuTile
ct.printf("value: %f\n", x)
ct.assert_(x > 0)

# Triton
tl.device_print("value", x)
tl.device_assert(x > 0, "x must be positive")

Atomic Operations

Synchronization

Host Functions

Data Types

Launch Patterns

cuTile (source) → Triton (target)

# cuTile
grid = ((N + BLOCK - 1) // BLOCK, 1, 1)
ct.launch(torch.cuda.current_stream(), grid, kernel, (x, y, N, BLOCK))

# Triton
grid = (triton.cdiv(N, BLOCK),)
kernel［grid］(x_ptr, y_ptr, N, BLOCK=256)

# cuTile 2D
grid = ((M + BLOCK_M - 1) // BLOCK_M, (N + BLOCK_N - 1) // BLOCK_N, 1)
ct.launch(stream, grid, kernel, (a, b, c, M, N, K, BLOCK_M, BLOCK_N))

# Triton 2D (tuple or lambda)
grid = (triton.cdiv(M, BLOCK_M), triton.cdiv(N, BLOCK_N))
kernel［grid］(a_ptr, b_ptr, c_ptr, M, N, K, BLOCK_M=64, BLOCK_N=64)

# Triton lambda (e.g. autotune)
grid = lambda meta: (triton.cdiv(M, meta['BLOCK_M']), triton.cdiv(N, meta['BLOCK_N']))
kernel［grid］(...)

cuTile → Triton Gotchas

1. Import: import cuda.tile as ct → import triton.language as tl (and add import triton if needed).
2. TMA for 2D+ loads (critical): cuTile uses TMA for block-aligned 2D+ loads. In Triton you must use tl.make_tensor_descriptor(...).load([...]) (see TensorDescriptor Pattern). Do not convert 2D+ ct.load(arr, index=, shape=) to raw tl.load(ptr + offs, mask=...) — that causes 500%-2000% (5-20x) regression.
3. Loop index → advance: Loop variable in ct.load(index=) → express as loop with tl.advance(ptr, delta) in Triton.
4. Gather/scatter → pointer: ct.gather/ct.scatter → tl.load/tl.store with ptr + offs and mask.
5. Type cast: ct.astype(x, dtype) → x.to(dtype).
6. Matrix multiply: ct.mma(a, b, acc=acc) → tl.dot(a, b, acc) (no acc= keyword).
7. Arange: ct.arange(N, dtype=ct.int32) → tl.arange(0, N).
8. Grid: Replace fixed 3-tuple and ct.launch(stream, grid, kernel, args) with grid = (...) or lambda meta: (...) and kernel［grid］(kernel_args).
9. None args: Dummy tensor + flag in cuTile → you can use None in Triton kernel args if the kernel supports it.
10. Host cdiv: (a + b - 1) // b → can use triton.cdiv(a, b) in Triton host.
11. Index math: Keep using Python +, *, // for indices; do not use tl.add/tl.mul for index arithmetic.
12. Kernel args: Tensor args in cuTile → pass pointers (and shapes/strides if needed) in Triton; constexpr/Constant → tl.constexpr.

Quick Reference Card (cuTile → Triton)

TensorDescriptor Pattern (ct.load/ct.store → Triton TMA)

Mandatory for 2D+ tile loads. When converting cuTile block-aligned loads (ct.load/ct.store with 2D+ shape) to Triton on Hopper (SM 90+) / Blackwell, use TensorDescriptor (tl.make_tensor_descriptor + .load/.store) so Triton can use TMA. Falling back to plain tl.load(ptr + offs, mask=...) for 2D+ block access causes 5-20x (500%-2000%) regression.

from triton.tools.tensor_descriptor import TensorDescriptor

def supports_host_descriptor():
    return torch.cuda.get_device_capability()[0] >= 9

@triton.jit
def _maybe_make_tensor_desc(desc_or_ptr, shape, strides, block_shape):
    if isinstance(desc_or_ptr, tl.tensor_descriptor):
        return desc_or_ptr
    else:
        return tl.make_tensor_descriptor(desc_or_ptr, shape, strides, block_shape)

@triton.jit
def kernel(X, X_shape_0, X_shape_1, X_stride_0, X_stride_1, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    X_desc = _maybe_make_tensor_desc(X, shape=[X_shape_0, X_shape_1],
                                      strides=[X_stride_0, X_stride_1],
                                      block_shape=[1, BLOCK])
    tile = X_desc.load([pid, 0])
    # ... computation ...
    X_desc.store([pid, 0], tile)

# Host
def wrapper(x):
    BLOCK = 128
    grid = (x.shape[0], 1, 1)
    if supports_host_descriptor():
        desc_x = TensorDescriptor(x, shape=x.shape, strides=x.stride(), block_shape=[1, BLOCK])
    else:
        desc_x = x
    kernel［grid］(desc_x, x.shape[0], x.shape[1], x.stride(0), x.stride(1), BLOCK)

• Shape/stride passed as individual scalars to the kernel.
• Use _maybe_make_tensor_desc for fallback when TensorDescriptor is not available.

Multi-dimensional Indexing

Array.slice → Triton (Ragged Tensors)

For Array.slice ragged tensor patterns, apply the mapping above: compute ptr + start * stride + offs manually and use tl.cdiv for tile counts.

ct.gather().item() → Triton (Runtime Index TMA)

For paged attention TMA: load the page index with tl.load(block_tables + idx), then use it as an offset into the cache tensor via tl.make_tensor_descriptor or manual pointer arithmetic.