xai-org/x-algorithm · error · ValueError

Unsupported swizzle triple for UMMA smem descriptor

Error message

Unsupported swizzle triple for UMMA smem descriptor

What it means

_layout_type maps a cute.Swizzle<B,M,S> triple to one of the UMMA shared-memory descriptor layout modes (SWIZZLE_NONE/32B/64B/128B/128B_BASE32B). Only M==4 with S==3 (canonical CUDA swizzles) or the exact triple Swizzle<2,5,2> are recognized; any other combination has no hardware encoding in this builder, so it rejects it. The error surfaces from make_smem_desc_base, which is the entry point used by all gemm_ptx kernels and smem_desc_base_from_tensor.

Source

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

def _layout_type(swizzle: cute.Swizzle) -> LayoutType:
    B, M, S = swizzle.num_bits, swizzle.num_base, swizzle.num_shift

    if M == 4:
        if S != 3:
            raise ValueError("Unexpected swizzle shift – want S==3 for M==4")
        return {
            0: LayoutType.SWIZZLE_NONE,
            1: LayoutType.SWIZZLE_32B,
            2: LayoutType.SWIZZLE_64B,
            3: LayoutType.SWIZZLE_128B,
        }[B]
    if M == 5:
        if (B, S) != (2, 2):
            raise ValueError("Only Swizzle<2,5,2> supported for 128B_BASE32B")
        return LayoutType.SWIZZLE_128B_BASE32B

    raise ValueError("Unsupported swizzle triple for UMMA smem descriptor")


def make_smem_desc_base(layout: cute.Layout, swizzle: cute.Swizzle, major: Major) -> int:
    layout_type = _layout_type(swizzle)

    VERSION = 1
    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:

View on GitHub (pinned to 24c60942c5)

Solutions

  1. Use a canonical swizzle: Swizzle<0,4,3> (none), <1,4,3> (32B), <2,4,3> (64B), <3,4,3> (128B), or Swizzle<2,5,2> for 128B_BASE32B
  2. Build the smem layout from a tile_to_shape over a standard Swizzle atom (e.g. cute.make_layout of tile_to_shape(SmemLayoutAtom, smem_shape)) instead of composing a custom swizzle
  3. If you need an unsupported swizzle mode, extend the LayoutType enum and the swizzle_atom_mn_size table in this file rather than passing the odd swizzle
  4. Print swizzle.num_bits/num_base/num_shift right before the call to confirm which triple you are actually generating

Example fix

# before
desc = make_smem_desc_base(layout, cute.Swizzle(3,3,3), Major.MN)
# after
desc = make_smem_desc_base(layout, cute.Swizzle(3,4,3), Major.MN)
Defensive patterns

Strategy: validation

Validate before calling

SUPPORTED = {(0,4,3),(1,4,3),(2,4,3),(3,4,3),(2,5,2)}
def is_supported_swizzle(sw):
    return (sw.num_bits, sw.num_base, sw.num_shift) in SUPPORTED

Type guard

def is_supported_swizzle(sw: cute.Swizzle) -> bool:
    return (sw.num_bits, sw.num_base, sw.num_shift) in {(0,4,3),(1,4,3),(2,4,3),(3,4,3),(2,5,2)}

Try / catch

try:
    desc = make_smem_desc_base(layout, sw, major)
except ValueError as e:
    if 'Unsupported swizzle' in str(e):
        sw = cute.Swizzle(3,4,3)  # fall back to 128B canonical
        desc = make_smem_desc_base(layout, sw, major)
    else:
        raise

Prevention

When it happens

Trigger: Calling make_smem_desc_base / smem_desc_base_from_tensor / gemm_ptx* with a Swizzle whose num_base (M) is neither 4 nor 5, e.g. Swizzle<3,3,3>; or M==4 with B>3 or S!=3 variants that escape the earlier check; or M==5 with (B,S)!=(2,2) (covered by the adjacent BASE32B error before reaching this line).

Common situations: Hand-rolled smem atom layouts copied from a different CUTLASS/CuTe DSL tutorial with a non-standard swizzle; porting Blackwell tcgen05 kernels that use sub-64B or 32B-base swizzle variants not supported here; upgrading the CUTLASS Python DSL where swizzle atom conventions changed.

Related errors

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