xai-org/x-algorithm · 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

K-major counterpart of error 371: after logical_divide(layout,(8,2)), the result must be congruent to ((1,1),(1,1)) — a strict hierarchical layout with an 8x2 atom and uniform strides above it. Any interleaving or non-divisible K extent (K not a multiple of 2) breaks the profile and is rejected before descriptor bit packing.

Source

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

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

    desc = 0
    desc |= (leading_byte_offset & 0x3FFF) << 16
    desc |= (stride_byte_offset & 0x3FFF) << 32
    desc |= (VERSION & 0x3) << 46
    desc |= (BASE_OFFSET & 0x7) << 49
    desc |= (LBO_MODE & 0x1) << 52
    desc |= (int(layout_type) & 0x7) << 61

    return desc & 0xFFFF_FFFF_FFFF_FFFF

View on GitHub (pinned to 24c60942c5)

Solutions

1. Ensure K is even and the layout is a clean tile_to_shape of an (8,2) K-major atom
2. Remove manual padding/offsets; let the swizzle handle bank conflicts
3. If a non-canonical layout is intentional, restructure it as ((8,m),(2,k)) with hierarchical strides before calling

Example fix

# before
bad = cute.make_layout((8,3), stride=(1,8))  # K=3 odd
# after
good = cute.make_layout((8,4), stride=(1,8))

assert cute.size(layout.shape[1]) % 2 == 0
assert cute.is_congruent(cute.logical_divide(layout,(8,2)), ((1,1),(1,1)))

When it happens

Trigger: K-major layout with K extent not divisible by 2, mixed stride hierarchy after dividing by (8,2), or a layout composed with a non-affine value layout so the divide leaves extra modes.

Common situations: Odd K sizes (K=1 or K=3 test cases); smem layouts with bank-conflict padding inserted manually that breaks the (8,2) tiling; layouts from cute.composition with swizzle applied before logical_divide.

Related errors

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

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