sgl-project/sglang · error · ValueError

Not a canonical UMMA_K Layout: Expected profile failure.

Error message

Not a canonical UMMA_K Layout: Expected profile failure.

What it means

For Major-K layouts, make_smem_desc_base splits the layout into (8, 2) atoms and requires the result be congruent to ((1,1),(1,1)) — a clean two-level hierarchy. If cute.is_congruent fails, the layout profile is not canonical and cannot be encoded into the smem descriptor.

Source

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

            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."
            )
        canonical_layout = cute.logical_divide(layout, (8, 2))
        if not cute.is_congruent(canonical_layout, ((1, 1), (1, 1))):
            raise ValueError("Not a canonical UMMA_K Layout: Expected profile failure.")
        stride_00 = canonical_layout.stride[0][0]
        if stride_00 != swizzle_atom_mn_size:
            raise ValueError("Not a canonical UMMA_K Layout: Expected stride failure.")
        stride_10 = canonical_layout.stride[1][0]
        if layout_type is not LayoutType.SWIZZLE_NONE and stride_10 != 1:
            raise ValueError("Not a canonical UMMA_K Layout: Expected stride failure.")
        stride_01 = canonical_layout.stride[0][1]
        stride_byte_offset, leading_byte_offset = stride_01, stride_10

    # ------------------------------------------------------------------ pack
    desc = 0
    # leading_byte_offset_  [16:30)
    desc |= (leading_byte_offset & 0x3FFF) << 16
    # stride_byte_offset_   [32:46)
    desc |= (stride_byte_offset & 0x3FFF) << 32
    # version_             [46:48)
    desc |= (VERSION & 0x3) << 46
    # base_offset_         [49:52)

View on GitHub (pinned to 0132848349)

Solutions

1. Ensure both layout extents are exact multiples of the (8, 2) atom split; pad if needed.
2. Rebuild the layout with cute.make_layout / cute.tile_to_shape so no remainder modes survive.
3. Test cute.is_congruent(cute.logical_divide(layout, (8, 2)), ((1,1),(1,1))) in a unit test before kernel launch.

canon = cute.logical_divide(layout, (8, 2))
assert cute.is_congruent(canon, ((1, 1), (1, 1)))

When it happens

Trigger: A Major.K layout with remainders in either mode, nested sub-modes, or non-uniform strides passed to gemm_ptx* / smem_desc_base_from_tensor.

Common situations: Building K-major smem tiles with cute.composition of mismatched atoms; layouts with a leftover mode from tile_to_shape when shape isn't a multiple of the atom.

Related errors

• Not a canonical UMMA_MN 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/657b745795c750ef. Report an issue: GitHub.