sgl-project/sglang · error · ValueError

M must be 64, 128 or 256

Error message

M must be 64, 128 or 256

What it means

make_instr_desc encodes the UMMA instruction M dimension into a 5-bit descriptor field that only represents M values of 64, 128 and 256. Any other M raises this ValueError before the bit-packing step.

Source

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

    a_neg: ScaleIn = ScaleIn.One,
    b_neg: ScaleIn = ScaleIn.One,
    c_sat: Saturate = Saturate.False_,
    is_sparse: bool = False,
    max_shift: MaxShift = MaxShift.NoShift,
) -> int:
    """
    Build the 32-bit instruction descriptor for Blackwell MMA.
    All matrix/accumulator **types must be CUTLASS scalar classes** –
    passing integers is forbidden.
    """
    # --- encode element formats -------------------------------------------------
    a_fmt = int(to_UMMA_format(a_type))
    b_fmt = int(to_UMMA_format(b_type))
    c_fmt = int(to_C_format(c_type))

    # --- range checks on M/N -----------------------------------------------------
    if M not in (64, 128, 256):
        raise ValueError("M must be 64, 128 or 256")
    if N < 8 or N > 256 or (N & 7):
        raise ValueError("N must be a multiple of 8 in the range 8…256")

    m_dim = M >> 4  # 5-bit field
    n_dim = N >> 3  # 6-bit field

    # fmt: off
    # --- pack the bit-fields -----------------------------------------------------
    desc = 0
    desc |= (0                 & 0x3) << 0        # sparse_id2 (always 0 here)
    desc |= (int(is_sparse)    & 0x1) << 2        # sparse_flag
    desc |= (int(c_sat)        & 0x1) << 3        # saturate
    desc |= (c_fmt             & 0x3) << 4        # c_format
    desc |= (a_fmt             & 0x7) << 7        # a_format
    desc |= (b_fmt             & 0x7) << 10       # b_format
    desc |= (int(a_neg)        & 0x1) << 13       # a_negate
    desc |= (int(b_neg)        & 0x1) << 14       # b_negate
    desc |= (int(a_major)      & 0x1) << 15       # a_major

View on GitHub (pinned to 0132848349)

Solutions

1. Snap your tile M to 64, 128, or 256 (e.g. use 128 as the default atom).
2. If a smaller logical M is needed, use a larger atom and mask/predicate the unused rows.
3. Check the caller (mma_op_to_idesc) to confirm which dimension of your instruction maps to descriptor M.

assert M in (64, 128, 256), f"M={M} is not a legal UMMA M"

When it happens

Trigger: Calling make_instr_desc(M=32, ...) or M=96/192 or any value outside (64, 128, 256); typically when deriving M from a tile shape that is not a legal UMMA atom size.

Common situations: Porting SM90 tile configurations to SM100; deriving the MMA M from an arbitrary GEMM tile_m without snapping it to a legal UMMA atom.

Related errors

• N must be a multiple of 8 in the range 8…256
• Unsupported tcgen05 MMA op kind: {type(op).__name__}
• Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}
• Unsupported CUTLASS scalar type for accumulator: {cutlass_ty
• Unexpected swizzle shift – want S==3 for M==4

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