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.

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.
Defensive patterns

Strategy: validation

Validate before calling

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

Prevention

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

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

Appendix: 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)