sgl-project/sglang · error · ValueError

Not a canonical UMMA_MN Layout: Expected stride failure.

Error message

Not a canonical UMMA_MN Layout: Expected stride failure.

What it means

After splitting a Major-MN layout into swizzle atoms, make_smem_desc_base requires the inner MN stride (stride_00) to be exactly 1 whenever the layout is swizzled. A non-unit inner stride means the layout is not contiguous within the atom and cannot be encoded.

Source

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

        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:
            raise ValueError("SWIZZLE_128B_BASE32B is invalid for Major-K")
        if not cute.size(layout.shape[0]) % 8 == 0:
            raise ValueError(
                "Not a canonical UMMA_K Layout: Expected MN-size multiple of 8."
            )

View on GitHub (pinned to 0132848349)

Solutions

1. Ensure the MN mode inside each atom is stride-1 contiguous (build with cute.make_layout using compact order).
2. Double-check the Major argument matches the actual layout construction.
3. Reuse the layout constructors from the shipped SM100 kernels.

canon = cute.logical_divide(layout, (atom_mn, atom_k))
if layout_type is not LayoutType.SWIZZLE_NONE:
    assert canon.stride[0][0] == 1

When it happens

Trigger: A Major.MN swizzled layout where the fastest-varying MN element inside each (atom_mn, atom_k) tile has stride != 1 — e.g. a column-major-ish MN mode or a layout with an interleaved stride pattern.

Common situations: Passing a K-major-style layout tagged as Major.MN, or constructing the layout with the mode order swapped.

Related errors

• Not a canonical UMMA_K Layout: Expected stride failure.
• Not a canonical UMMA_MN Layout: Expected profile failure.
• SWIZZLE_128B_BASE32B is invalid for Major-K
• Not a canonical UMMA_K Layout: Expected MN-size multiple of
• Not a canonical UMMA_K Layout: Expected profile failure.

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