sgl-project/sglang · error · ValueError

Unsupported swizzle triple for UMMA smem descriptor

Error message

Unsupported swizzle triple for UMMA smem descriptor

What it means

_layout_type raises when the swizzle triple (num_base, num_bits, num_shift) matches neither the M==4 (Swizzle<*,4,3>) nor M==5 (Swizzle<2,5,2>) families — i.e. it is not a UMMA-legal shared-memory layout at all.

Source

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

    if M == 4:  # Swizzle<*,4,3>
        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
        ]  # KeyError ⇒ invalid B→ raise
    if M == 5:  # Swizzle<2,5,2> (the only legal triple for M==5)
        if (B, S) != (2, 2):
            raise ValueError("Only Swizzle<2,5,2> supported for 128B_BASE32B")
        return LayoutType.SWIZZLE_128B_BASE32B

    # Any other (M,B,S) triple is not a UMMA-legal shared-memory layout
    raise ValueError("Unsupported swizzle triple for UMMA smem descriptor")


def make_smem_desc_base(
    layout: cute.Layout, swizzle: cute.Swizzle, major: Major
) -> int:
    """
    Convert a 2-D *shared-memory* Cute layout into the Blackwell 64-bit
    smem-descriptor, without the smem start address.
    layout must correspond to layout of an uint128 tensor.
    """
    # ------------------------------------------------------------------ meta
    layout_type = _layout_type(swizzle)  # resolve SWIZZLE_* family

    VERSION = 1  # bits 46–47
    LBO_MODE = 0  # bit  52
    BASE_OFFSET = 0  # bits 49–51   (CUTLASS always 0)

    # ---------------------------------------------------------- strides  (units: uint128_t = 16 B)

View on GitHub (pinned to 0132848349)

Solutions

1. Use LayoutType.SWIZZLE_NONE-compatible layouts (no swizzle) or one of the legal triples: Swizzle<0..3,4,3> or Swizzle<2,5,2>.
2. Copy the swizzle from a working kernel's atom in this same directory rather than computing one.
3. Consult the UMMA smem descriptor table (32B/64B/128B swizzle atoms) when choosing layouts.

def is_umma_legal_swizzle(sw):
    B, M, S = sw.num_bits, sw.num_base, sw.num_shift
    return (M == 4 and S == 3 and 0 <= B <= 3) or (M == 5 and (B, S) == (2, 2))

When it happens

Trigger: Calling make_smem_desc_base with arbitrary swizzles such as Swizzle<1,3,0>, Swizzle<3,4,3>, or a non-swizzled layout's default swizzle object with unusual fields.

Common situations: Deriving swizzle parameters programmatically from tensor shapes/strides and landing on an unsupported combination; using SM90 wgmma-era swizzles with the SM100 UMMA descriptor builder.

Related errors

• Unexpected swizzle shift – want S==3 for M==4
• Only Swizzle<2,5,2> supported for 128B_BASE32B
• Not a canonical UMMA_MN Layout: Expected profile failure.
• SWIZZLE_128B_BASE32B is invalid for Major-K
• Not a canonical UMMA_K Layout: Expected profile failure.

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