Gotchas — Most Common cuTile → Triton Translation Errors
Comprehensive table of patterns that frequently break or regress when porting
@ct.kernel to @triton.jit. Read this BEFORE writing the Triton kernel — most
entries describe a wrong-by-default first attempt.
| Pattern | cuTile | Triton | Common Mistake |
|---|---|---|---|
| mma accumulator | ct.mma(a, b, acc=acc) |
tl.dot(a, b, acc) |
Using keyword acc= in Triton (positional only) |
| mma float32→tf32 | Explicit ct.astype(..., ct.tfloat32) guard before ct.mma |
tl.dot(a, b, allow_tf32=True) (default) |
Over-specifying; Triton auto-casts by default |
| Type cast | ct.astype(x, dtype) |
x.to(dtype) |
Using ct.astype in Triton |
| Grid | (n, 1, 1) tuple, ct.launch(stream, grid, kernel, args) |
lambda meta: (n,) or tuple, bracket launch |
Using ct.launch or 3-tuple in Triton |
| Host cdiv | (a + b - 1) // b (Python) |
triton.cdiv(a, b) |
Forgetting triton.cdiv in host |
| 2D+ tile load | ct.load(arr, index=(i,j), shape=(BM,BK)) (cuTile uses TMA) |
tl.make_tensor_descriptor(...).load([...]) |
Using raw tl.load(ptr+offs, mask=m) → 5-20x regression; always use TMA for 2D+ block loads |
| Index type | Block index in ct.load/ct.store | Element offset (ptr + offs) or TMA descriptor | Using block index as tl.load offset |
| arange | ct.arange(N, dtype=ct.int32) |
tl.arange(0, N) |
Triton has start param (0, N) |
| None args | Dummy tensor + flag | Allowed in kernel | Carrying over dummy+flag when not needed |
| String const | ct.Constant[int] only (no str) |
tl.constexpr (any type) |
Keeping int enum; Triton can use str constexpr if needed |
| Shape args | Static/constexpr in ct.full/ct.zeros | Dynamic shapes OK in Triton | Over-constraining shapes |
| Launch | ct.launch(stream, grid, kernel, args) |
bracket launch (grid then args) | Leaving ct.launch in Triton host |
| Branch vars | Pre-define before if | Can define in branch | Over-defining before branch in Triton |
| Pointer table type | Typed tensor descriptor (auto) | tl.load(ptrs+idx).to(tl.pointer_type(DTYPE)) where DTYPE: tl.constexpr |
Hardcoding tl.float16 → cudaErrorIllegalAddress for bfloat16/float32 inputs |
| Stride dtype | cuTile uses tensor shape (auto) | Pass strides as torch.int64, not int32 |
int32 overflows → illegal address for large matrices (M×K > 2^31) |
| dtype map coverage | cuTile typed tensors (auto) | _DTYPE_MAP must cover all dtypes (incl. float8); use hasattr guards |
Missing entry → ValueError: Unsupported dtype before kernel launch |
| tl.math.erf dtype | cuTile erf handles all dtypes | tl.math.erf only accepts fp32/fp64 |
ValueError: Expected dtype ['fp32', 'fp64'] but got fp16 — do NOT replace with tanh approximation (mathematically wrong); let Triton auto-promote or cast input |
| tl.exp with fp16 | cuTile exp handles all dtypes | Cast to fp32 before tl.exp for precision: tl.exp(x.to(tl.float32)) |
Precision loss or NaN with fp16 inputs in exp/log/sqrt |
| Math func approx | N/A | Never substitute tl.math.erf with tanh-based approximation |
Using GELU tanh formula (0.044715*x³) as erf approximation is mathematically incorrect — they are different functions |
Layout flag (transpose) |
cuTile may use one path per layout | Need two Triton kernels when math differs (e.g. MLA: qk [H,N] vs [N,H], different V TMA) |
Reusing transpose-only logic for transpose=False + fixed blocks → 3–15× on that mode; see ../translations/advanced-patterns.md |
| Batched matmul | ct.matmul(W, X) broadcasts implicitly at tile level |
tl.dot(W, X) only supports 2D operands |
Using broadcast_to + tl.dot → 10-50× slower, no tensor cores (see FFT anti-pattern in performance-gotchas.md) |
| Batch-per-block | cuTile processes 1 batch per block naturally | Triton temptation: process BS batches per block | Creates BS× register pressure, breaks tensor core compatibility |