xai-org/x-algorithm · 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 the congruence check passes, make_smem_desc_base verifies that the stride of the innermost MN mode of the swizzle atom is 1 when a swizzle is active. A swizzled MN-major atom must be contiguous in the MN dimension (stride_00 == 1) so the swizzle function can permute the low address bits; anything else cannot be encoded in the descriptor's stride fields.

Source

Thrown at phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py:251

    LBO_MODE = 0
    BASE_OFFSET = 0

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

View on GitHub (pinned to 24c60942c5)

Solutions

  1. Make the innermost MN stride 1: build the atom as make_layout((atom_mn,atom_k), stride=(1,atom_k)) before tile_to_shape
  2. If the operand really is K-major, pass Major.K instead of Major.MN
  3. If you genuinely want non-contiguous MN, use SWIZZLE_NONE (Swizzle<0,4,3>) which skips this check

Example fix

# before
atom = cute.make_layout((8,8), stride=(8,1))
# after
atom = cute.make_layout((8,8), stride=(1,8))
Defensive patterns

Strategy: validation

Validate before calling

cl = cute.logical_divide(layout, (atom_mn, atom_k))
if layout_type != LayoutType.SWIZZLE_NONE:
    assert cl.stride[0][0] == 1, 'inner MN stride must be 1 for swizzled MN-major layouts'

Prevention

When it happens

Trigger: MN-major layout with a swizzle (SWIZZLE_32B/64B/128B/BASE32B) where the innermost MN stride is not 1 — e.g. a transposed atom where the fastest-varying mode is K, or an atom built with stride (atom_k, 1) instead of (1, atom_k).

Common situations: Swapping M/N-major and K-major conventions when porting a kernel; reusing a K-major smem layout while declaring Major.MN; layouts derived from coalesce() or composition that changed the inner stride order.

Related errors

AI-assisted analysis of xai-org/x-algorithm@24c60942c5 (2026-08-28). Data as JSON: /api/errors/261983037b946e95. Report an issue: GitHub.