sgl-project/sglang · error · ValueError

Not a canonical UMMA_MN Layout: Expected profile failure.

Error message

Not a canonical UMMA_MN Layout: Expected profile failure.

What it means

make_smem_desc_base validates that a Major-MN shared-memory layout is congruent to the canonical swizzle-atom tiling ((atom_mn, atom_k) split into ((1,1),(1,1))). If cute.is_congruent fails, the layout does not have the profile the descriptor encoder assumes, so it refuses rather than encoding garbage.

Source

Thrown at python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py:254

    LBO_MODE = 0  # bit  52
    BASE_OFFSET = 0  # bits 49–51   (CUTLASS always 0)

    # ---------------------------------------------------------- strides  (units: uint128_t = 16 B)
    swizzle_atom_mn_size = {
        LayoutType.SWIZZLE_NONE: 1,
        LayoutType.SWIZZLE_32B: 2,
        LayoutType.SWIZZLE_64B: 4,
        LayoutType.SWIZZLE_128B: 8,
        LayoutType.SWIZZLE_128B_BASE32B: 8,
    }[layout_type]

    if major is Major.MN:
        swizzle_atom_k_size = 4 if layout_type is LayoutType.SWIZZLE_128B_BASE32B else 8
        canonical_layout = cute.logical_divide(
            layout, (swizzle_atom_mn_size, swizzle_atom_k_size)
        )
        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):
            raise ValueError(
                "Not a canonical UMMA_MN Layout: Expected profile failure."
            )
        stride_00 = canonical_layout.stride[0][0]
        if layout_type is not LayoutType.SWIZZLE_NONE and stride_00 != 1:
            raise ValueError("Not a canonical UMMA_MN Layout: Expected stride failure.")
        stride_10 = canonical_layout.stride[1][0]
        if stride_10 != swizzle_atom_mn_size:
            raise ValueError("Not a canonical UMMA_MN Layout: Expected stride failure.")
        stride_01, stride_11 = (
            canonical_layout.stride[0][1],
            canonical_layout.stride[1][1],
        )
        if layout_type is LayoutType.SWIZZLE_NONE:
            stride_byte_offset, leading_byte_offset = stride_01, stride_11
        else:
            stride_byte_offset, leading_byte_offset = stride_11, stride_01
    else:
        if layout_type == LayoutType.SWIZZLE_128B_BASE32B:

View on GitHub (pinned to 0132848349)

Solutions

1. Construct the layout with cute.make_layout / cute.tile_to_shape so it is an exact product of the swizzle atom (atom_mn x 8, or 4 for 128B_BASE32B).
2. Compare your layout against the ones built inside the shipped gemm_ptx* helpers and match their construction path.
3. Falling back to smem_desc_base_from_tensor on a tensor whose cute view already has a canonical layout avoids hand-rolled profiles.

import cute

atom_k = 4 if is_base32b else 8
canon = cute.logical_divide(layout, (atom_mn, atom_k))
assert cute.is_congruent(canon, ((1, 1), (1, 1)))

When it happens

Trigger: Passing a Major.MN layout whose shape/stride structure is not a clean product of (swizzle_atom_mn_size, swizzle_atom_k_size) tiles — e.g. with remainders, nested hierarchies, or stride holes — via gemm_ptx* or smem_desc_base_from_tensor.

Common situations: Building smem tile layouts from tensor views with padding or non-uniform strides; composing layouts with cute.tile_to_shape incorrectly so the K or MN remainder leaks into the profile.

Related errors

• Not a canonical UMMA_K Layout: Expected profile failure.
• Unexpected swizzle shift – want S==3 for M==4
• Only Swizzle<2,5,2> supported for 128B_BASE32B
• Unsupported swizzle triple for UMMA smem descriptor
• Not a canonical UMMA_MN Layout: Expected stride failure.

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/278103e8d621ae2b. Report an issue: GitHub.