4.1.1 Branch Predication Bugs

Branch predication bugs arise when control flow expressed over logical tiles is incorrectly lowered to operate over physical threads. Incorrect lowering of predicates can cause valid tile instances to be skipped or incorrectly masked.

# Simplified from the actual issue #5265 reproducer
@triton.jit
def repro_kernel(q_ref, k_ref, v_ref, output_ptr):
    offsets64  = tl.arange(0, 64)
    offsets128 = tl.arange(0, 128)

    q  = tl.load(q_ref + (offsets64[:, None] * 128 + offsets128[None, :]))   # [64, 128]
    k  = tl.load(k_ref + (offsets128[:, None] * 64 + offsets64[None, :]))    # [128, 64]
    qk = tl.dot(q, k).to(tl.bfloat16)                                       # [64, 64]

    v  = tl.load(v_ref + (offsets64[:, None] * 128 + offsets128[None, :]))   # [64, 128]
    o  = tl.dot(qk, v)                                                       # [64, 128]

    tl.store(output_ptr + (offsets64[:, None] * 128 + offsets128[None, :]), o.to(tl.bfloat16))

# ✅ num_warps=4 — compiles and executes correctly
# ❌ num_warps=8 — triggers assertion failure in LinearLayout.cpp::reshapeOuts(...)
repro_kernel[(1, 1)](q, k, v, output, num_warps=8, num_ctas=1, num_stages=3)

The failure is a warp-configuration-sensitive regression in Triton's Linear Layout pass. The commit that introduced stmatrix support through linear layouts (49266aa) contains assumptions about tile-to-warp mapping that hold for num_warps=4 but break for num_warps=8. Because the Linear Layout pass is responsible for deriving how logical tiles are distributed across physical warps and threads, including the predication and masking decisions that depend on this mapping, the invalid layout causes the compiler to fail before correct control flow can be generated.

4.1.2 Instruction Scheduling Bugs

Instruction scheduling bugs arise when backend reordering violates compiler-synthesized tile pipeline invariants, rather than explicit source-level data dependencies.

# Schematic illustration (not actual code from the issue)

# AMD backend's createLocalPrefetchSchedule() assumes this IR order:
# ┌──────────────────────────────────────────────────┐
# │  local_stores   ← write prev-iteration data to LDS│
# │  global_loads   ← prefetch next tile from GMEM    │
# │  compute        ← process current tile            │
# │  local_loads    ← read prefetched data from LDS   │
# └──────────────────────────────────────────────────┘

# ❌ After ReorderInstructionsPass::scheduleGlobalLoadLocalStore():
# ┌──────────────────────────────────────────────────┐
# │  global_loads   ← moved ahead of local_stores!    │
# │  local_stores   ← moved behind global_loads!      │
# │  compute / local_loads                            │
# └──────────────────────────────────────────────────┘
# Additionally, sync/barrier before local_stores causes
# sched.group masks to no longer work as intended.

The reordering pass and the prefetch scheduling pass operate under conflicting assumptions about IR ordering. The reorder pass rearranges instructions without preserving the ordering invariant that createLocalPrefetchSchedule depends on.

4.1.3 Warp Control Bugs

Warp control bugs arise from incorrect synthesis or propagation of warp-level execution metadata across compiler stages.

# Schematic illustration (specific values are illustrative)

# TritonGPU IR invariant (as described in the issue):
# For any tensor with a #blocked layout where warpsPerCta = [x, y]:
#     x * y  MUST equal  module's triton_gpu.num-warps

# WSMaterialization pass modifies triton_gpu.num-warps
# but does NOT modify tensor layouts within the module.
#
# → warpsPerCta fields in tensor layouts become stale
# → The invariant (x * y == num-warps) is broken
# → The resulting IR is invalid
#
# Note: the issue author points out that simply updating warpsPerCta
# may NOT be the correct fix—a new layout type may be needed.

WSMaterialization performs a partial update, modifying num-warps without synchronizing warpsPerCta fields. This breaks a layout invariant that was previously unchecked by the IR verifier. The issue author suggests the representation may need to be extended to express warp specialization groups.

4.2.1 IR Construction Bugs

IR construction bugs arise during front-end lowering where non-determinism or inconsistency in the creation of tile-level IR entities produces semantically unstable IRs.

# TileLang FlashAttention backward kernel (func.script() dump)
            # T.handle ordering before T.block_attr appears to be random
            
            # ─── Compilation Run 1 ───
            dv_shared = T.handle("float16", "shared.dyn")
            K_shared  = T.handle("float16", "shared.dyn")
            dk_shared = T.handle("float16", "shared.dyn")
            T.block_attr({"layout_map": {
                dQ.data:   metadata["tl.Layout"][0],
                dv_shared: metadata["tl.Layout"][1],
                K_shared:  metadata["tl.Layout"][2],
                dk_shared: metadata["tl.Layout"][3]}})
            
            # ─── Compilation Run 2 (same source code!) ───
            K_shared  = T.handle("float16", "shared.dyn")   # ← different order!
            dv_shared = T.handle("float16", "shared.dyn")   # ← different order!
            dk_shared = T.handle("float16", "shared.dyn")
            T.block_attr({"layout_map": {
                K_shared:  metadata["tl.Layout"][0],         # ← bindings shifted!
                dQ.data:   metadata["tl.Layout"][1],
                dv_shared: metadata["tl.Layout"][2],
                dk_shared: metadata["tl.Layout"][3]}})
            
            # → tilelang.compile fails to load cached kernel
            # → logically identical kernels treated as distinct artifacts

Non-deterministic creation order of shared-memory tile handles causes layout_map bindings to differ across compilations, defeating kernel caching. The compiler lacks a canonical ordering step for tile handle construction.

4.2.2 IR Transformation Bugs

IR transformation bugs arise when mid-end optimization passes rewrite tile IRs in ways that invalidate tile-specific execution invariants, such as boundary conditions and masked regions.

// Minimal reproducer from the issue
            #include <Halide.h>
            using namespace Halide;
            
            int main() {
                Var x{"x"};
                Func f{"f"}, g{"g"};
            
                f(x) = select(x == 0, x + 1, undef<int>());
                g(x) = select(x == 0, f(x), -f(1 - x));
            
                g.trace_stores();
            
                Buffer<int> output = g.realize({4});
                // Expected: g(0)=1, g(1)=-1, ...
                // Actual output:
                //   Store g.0(1) = -1
                //   g(0) is never computed  ← ❌ BUG
            }

The optimizer exploits undef values to prune a branch of select, inadvertently eliminating the x == 0 case where the computation is well-defined. This causes the compiler to silently reduce the effective iteration domain, skipping points that the original tiled logic correctly covers.

4.3.1 Launch Configuration Bugs

Launch configuration bugs arise when the compiler miscomputes launch parameters derived from tile-level abstractions.

# Reproducer from the issue
            import torch, functools
            from torch._inductor.utils import run_and_get_triton_code
            from torch._inductor import config
            
            config.triton.max_tiles = 3
            config.triton.prefer_nd_tiling = True
            
            full_size, view_size = (5, 5, 5, 5, 5), (3, 3, 5, 3, 5)
            
            def get_input():
                full = torch.randn(full_size).to("cuda")
                return torch.as_strided(full, view_size, full.stride())
            
            a, b = get_input(), get_input()
            opt_fn = torch.compile(functools.partial(torch.add))
            code = run_and_get_triton_code(opt_fn, a, b)
            
            # ❌ BackendCompilerFailed / AssertionError
            # The LoopBody iteration prefix "z" conflicts with the
            # z-dimension range tree prefix (banned in the symbol system).
            # No corresponding ZBLOCK is defined.
            # → Internal symbol-prefix collision during code generation

The codegen infrastructure assumes iteration prefixes for tile groups won't collide with reserved dimension prefixes. Increasing the tile group count to 3 introduces the z prefix which overlaps with the reserved z-dimension namespace.

4.3.2 Thread-Block Mapping Bugs

Thread–block mapping bugs arise when the compiler emits incorrect logic for mapping logical tiles to physical grid and block indices.

# Reproducer from the issue
            import torch
            from torch.nn.attention.flex_attention import flex_attention
            
            flex_attention = torch.compile(flex_attention, dynamic=True)
            
            DEVICE = 'cuda:0'
            dtype = torch.float32
            
            B, H, L, C = (22720, 3, 64, 32)   # ❌ Error
            # B, H, L, C = (20480, 3, 64, 32) # ✅ No Error
            # B, H, L, C = (20480, 4, 64, 32) # ❌ Error
            
            q = torch.rand(B, H, L, C, dtype=dtype, device=DEVICE)
            k = torch.rand(B, H, L, C, dtype=dtype, device=DEVICE)
            v = torch.rand(B, H, L, C, dtype=dtype, device=DEVICE)
            
            attn = flex_attention(q, k, v)
            # RuntimeError: Triton Error [CUDA]: invalid argument
            # at: self.launch(gridX, gridY, gridZ, ...)

When batch and head dimensions are large enough, the compiler maps them onto grid axes whose product exceeds CUDA hardware launch limits. The tile-to-grid mapping logic does not check for overflow when flattening multi-dimensional tile iteration spaces into physical grid dimensions, causing the kernel launch to fail with an invalid argument error.

4.4.1 Indexing and Strides Bugs

Indexing and strides bugs arise when the compiler incorrectly synthesizes address computations from logical tile indices.

# Issue: "Segfault in dds_matmul"
# Reproducer triggers _dsd_kernel via blocksparse matmul backward pass:

from triton.ops.blocksparse.matmul import _matmul

# ctx contains pre-built LUT tensors for the sparse layout
# dc is the upstream gradient tensor
_matmul.backward(ctx=ctx, dc=dc)

# Under cuda-memcheck (PYTORCH_NO_CUDA_MEMORY_CACHING=1):
#   "Invalid __global__ read of size 4"
#   "    at 0x00000670 in _dsd_kernel"
#   "    by thread (97,0,0) in block (0,17,0)"
#   "    Address 0x7f1306db0400 is out of bounds"
#
# The segfault is rare but reliably reproduced under cuda-memcheck.
# Adding logging or capturing intermediate state perturbs timing
# enough to make the bug disappear (heisenbug behavior).

The _dsd_kernel (dense-sparse-dense matmul) performs out-of-bounds global memory reads at runtime due to incorrect address computation within tile-local index arithmetic for sparse access patterns. The issue author notes heisenbug behavior: perturbations to timing or memory layout suppress the fault.

4.4.2 Resource Allocation Bugs

Resource allocation bugs arise when the compiler incorrectly infers the lifetime or reuse of registers or shared-memory buffers across tile phases.

# Simplified from the full deepgemm reproducer in the issue
# Config: thread_num=256, warp specialization enabled, Hopper GPU

for k in T.Pipelined(T.ceildiv(K, block_K), num_stages=num_stages):
    T.copy(A[by * block_M, k * block_K], A_shared)
    T.copy(B[bx * block_N, k * block_K], B_shared)

    # Load scale into shared memory via SIMT copy
    for i in T.Parallel(block_M):
        Scale_C_shared[i] = scales_a[by * block_M + i, k] * scales_b[bx, k]

    T.gemm(A_shared, B_shared, C_local, transpose_B=True, policy=policy)

    for i, j in T.Parallel(block_M, block_N):
        C_local_accum[i, j] += C_local[i, j] * Scale_C_shared[i]

# ❌ Deadlocks on NVIDIA Hopper GPU with 8 warps
# pass_configs: tl.disable_warp_specialized = False

The issue reports a deadlock when running a deepgemm kernel with 8 warps and warp specialization on Hopper. Our analysis attributes this to premature register deallocation (warpgroup_reg_dealloc) during tile-parallel copy: the compiler frees staging registers in the global → register → shared path before the SIMT copy completes, corrupting shared memory values and causing the subsequent warp-synchronized operations to hang.

4.4.3 Ordering and Caching Bugs

Ordering and caching bugs arise from incorrect kernel cache reuse or missing synchronization. A common pattern is the cache's specialization key omitting semantically relevant parameters.

# Issue: "Kernels that use the Cholesky solver might fail
#         when block_dim changes"
#
# cuMathDx needs block_dim at compile time for Cholesky solver.
# Compiled result is stored as an LTO file.

# ─── First launch ───
wp.launch(cholesky_kernel, dim=N, block_dim=128)
# → Compiles LTO with block_dim=128, caches it ✅

# ─── Second launch (different block_dim) ───
wp.launch(cholesky_kernel, dim=N, block_dim=256)
# → Kernel hash does NOT include block_dim
# → Cache returns LTO compiled for block_dim=128
# → ❌ Incorrect behavior: wrong block_dim assumptions
#
# Note: the issue does not include reproducer code;
# the above is a schematic illustration of the reported behavior.

The kernel caching mechanism does not include block_dim in its hash key. When block_dim changes between launches, the cache returns an LTO file compiled with incorrect block dimension assumptions, leading to wrong results or failures.

4.5.1 Special-Value Handling Bugs

Special-value handling bugs occur when the compiler mishandles floating-point edge cases (NaN, denormals, signed zeros) during tile-level execution.

# Reproducer from the issue
import numpy as np
import warp as wp

def tile_cholesky(M, L, tilesize):
    @wp.kernel
    def cholesky(M: wp.array(dtype=wp.float32, ndim=2),
                 L: wp.array(dtype=wp.float32, ndim=2)):
        M_tile = wp.tile_load(M, shape=(tilesize, tilesize))
        L_tile = wp.tile_cholesky(M_tile)
        wp.tile_store(L, L_tile)
    wp.launch_tiled(cholesky, dim=(1,), inputs=[M, L], block_dim=32)

dim = 5;  tilesize = 4  # tile 4×4, but source array is 5×5
Mwp = wp.array(np.eye(dim) * 2.0, dtype=wp.float32)
Lwp = wp.zeros((dim, dim), dtype=wp.float32)
tile_cholesky(Mwp, Lwp, tilesize)

err = np.linalg.cholesky(np.eye(dim) * 2.0)[:tilesize, :tilesize] \
    - Lwp.numpy()[:tilesize, :tilesize]
# ❌ Unexpected results including NaNs

wp.tile_load does not correctly handle 2D source arrays with strides incompatible with the tile shape. When the source array is 5×5 but the tile is 4×4, the mismatch in memory layout causes some tile lanes to read uninitialized or out-of-bounds values, producing NaNs. The CHANGELOG fix: "Fix 2D tile load when source array and tile have incompatible strides."

4.5.2 Data-Type Semantics Bugs

Data-type semantics bugs arise when incorrect type propagation breaks the mathematical structure of tile-level layout and index transformations.

# TVM MetaSchedule — int8 conv2d workload
# apply_trace generated by print(sch.trace) fails.
# Error stack: TransformLayout → transform_block_layout
#            → IndexMap::NonSurjectiveInverse
#
# These transforms support tiled layout algebra:
#   (n, y, x) → n × 64 + y × 8 + x
# Schedule trace contains index maps using T.int64(...).
# Data-type mismatch with int8 workload breaks these transforms.
# → Compilation failure

Tiled layout transformation passes fail when applied to an int8 workload, with errors in TransformLayout and IndexMap::NonSurjectiveInverse. The precise internal mechanism is not diagnosed in the issue, but the pattern is data-type mismatch breaking tiled layout transforms.

4.5.3 Operator Implementation Bugs

Operator implementation bugs arise when the compiler synthesizes incorrect logic for operators under tiling, masking, or fusion.

# Reproducer from the issue (also Figure 3 in our paper)
@triton.jit
def matmul_argmax_kernel(A, B, Out,
    BLOCK_M: tl.constexpr, BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr):
    offs_m = tl.arange(0, BLOCK_M)
    offs_n = tl.arange(0, BLOCK_N)
    offs_d = tl.arange(0, BLOCK_K)

    a_ptrs = A + offs_m[:, None] * BLOCK_K + offs_d[None, :]
    b_ptrs = B + offs_n[:, None] * BLOCK_K + offs_d[None, :]

    a = tl.load(a_ptrs)
    b = tl.load(b_ptrs)

    dist = tl.dot(a, tl.trans(b))
    assignments = tl.argmax(dist, 1)
    tl.store(Out + offs_m, assignments)

M, N, K = 32, 32, 32
a = torch.rand(M, K, dtype=torch.float16, device="cuda")
b = torch.rand(N, K, dtype=torch.float16, device="cuda")
out = torch.zeros(M, dtype=torch.int32, device="cuda")
matmul_argmax_kernel[(1,)](a, b, out, BLOCK_M=M, BLOCK_N=N, BLOCK_K=K)
# ❌ SIGSEGV during compilation
# gdb: crash in mlir::OperationFolder::tryToFold()
#   within TritonGPURemoveLayoutConversionsPass

A semantically valid tl.dot + tl.argmax composition triggers a segfault during compilation. The gdb trace shows the crash occurs in TritonGPURemoveLayoutConversionsPass when the pattern rewriter attempts to fold operations produced by the argmax reduction lowering. In our paper's analysis, the compiler's synthesized reduction algorithm contains incorrect assumptions about the output layout after tl.dot.

# Simplified from the full reproducer in the issue
@T.prim_func
def main(A: T.Buffer((M, K), "float16"),
         B: T.Buffer((K, N), "float16"),
         C: T.Buffer((M, N), "float16")):
    with T.Kernel(T.ceildiv(N, 192), T.ceildiv(M, 64), threads=128*2) as (bx, by):
        A_shared = T.alloc_shared((64, 32), "float16")     # block_M=64, block_K=32
        B_shared = T.alloc_shared((32, 192), "float16")    # block_K=32, block_N=192
        C_local  = T.alloc_fragment((64, 192), "float")

        T.clear(C_local)
        for k in T.Pipelined(T.ceildiv(K, 32), num_stages=3):
            T.copy(A[by * 64, k * 32], A_shared)
            T.copy(B[k * 32, bx * 192], B_shared)
            T.gemm(A_shared, B_shared, C_local, policy=T.GemmWarpPolicy.FullCol)
        T.copy(C_local, C[by * 64, bx * 192])

# 2 warpgroups → per-warpgroup MMA: M64, N96, K32
# Reporter states this should be supported (multiple of wgmma.m64n16k16)
#
# ❌ Compilation fails on NVIDIA H100 with CUTLASS/CUTE error:
#    "Static shape_div failure"
#    at cute::shape_div(_96, _64) — 96 and 64 are not divisible

The per-warpgroup MMA shape (M64, N96, K32) cannot be evenly decomposed into the SM90 MMA atom layout by the CUTE backend. Specifically, shape_div(96, 64) fails because 96 is not divisible by 64. The compiler does not validate tile-to-MMA compatibility before committing to the lowering path, causing the failure to surface as a backend static assertion rather than an informative diagnostic.