sgl-project/sglang · error · ValueError

N must be a multiple of 8 in the range 8…256

Error message

N must be a multiple of 8 in the range 8…256

What it means

The UMMA instruction descriptor's N field is 6 bits storing N>>3, so N must be a multiple of 8 between 8 and 256 inclusive. make_instr_desc raises this ValueError when N is out of range or not 8-aligned.

Source

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

    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
    desc |= (int(b_major)      & 0x1) << 16       # b_major
    desc |= (n_dim             & 0x3F) << 17      # n_dim (6 bits)

View on GitHub (pinned to 0132848349)

Solutions

1. Round N up to the nearest multiple of 8 within [8, 256] and predicate the excess columns.
2. Verify head_dim/tile_k values are 8-aligned; typical values 64/128/256 are safe.
3. If N > 256 you need multiple MMA instructions, not a larger descriptor N.

assert 8 <= N <= 256 and N % 8 == 0, f"N={N} must be multiple of 8 in [8, 256]"

When it happens

Trigger: Calling make_instr_desc with N values like 4, 12, 300, or any non-multiple of 8; common when deriving N from a head_dim or tile N of e.g. 96 is fine but 100 or 264 is not.

Common situations: Using non-standard head dims or K-tile sizes when building SM100 attention/GEMM descriptors; porting shapes from architectures with different alignment rules.

Related errors

• M must be 64, 128 or 256
• Not a canonical UMMA_K Layout: Expected MN-size multiple of
• 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

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