sgl-project/sglang · error · TypeError

Unsupported CUTLASS scalar type for accumulator: {cutlass_ty

Error message

Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}

What it means

to_C_format maps the accumulator scalar type to the 2-bit encoding in the UMMA instruction descriptor and only supports Float16, Float32 and Int32. Any other CUTLASS scalar class raises this TypeError from make_instr_desc.

Source

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

    # Float-8 / Float-6 / Float-4 – add whenever CUTLASS exposes them
    if cutlass_type is cutlass.Float8E4M3FN:
        return MXF8F6F4Format.E4M3
    if cutlass_type is cutlass.Float8E5M2:
        return MXF8F6F4Format.E5M2
    raise TypeError(f"Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}")


def to_C_format(cutlass_type) -> int:
    """
    Map a CUTLASS scalar class to the 2-bit accumulator encoding.
    """
    if cutlass_type is cutlass.Float16:
        return CFormat.F16
    if cutlass_type is cutlass.Float32:
        return CFormat.F32
    if cutlass_type is cutlass.Int32:
        return CFormat.S32
    raise TypeError(
        f"Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}"
    )


# ---------------------------------------------------------------------------
# The constructor – accepts only CUTLASS scalar classes
# ---------------------------------------------------------------------------


def make_instr_desc(
    a_type,  # CUTLASS scalar class, e.g. cutlass.Int8
    b_type,
    c_type,
    M: int,  # 64, 128 or 256
    N: int,  # 8 … 256 (multiple of 8)
    a_major: Major,
    b_major: Major,
    a_neg: ScaleIn = ScaleIn.One,

View on GitHub (pinned to 0132848349)

Solutions

1. Use cutlass.Float32 for floating accumulation, cutlass.Float16, or cutlass.Int32.
2. For BF16 compute, accumulate in Float32 and convert — do not request a BF16 C format.
3. Verify the argument order of make_instr_desc so c_type is actually the accumulator.

import cutlass

SUPPORTED_C = {cutlass.Float16, cutlass.Float32, cutlass.Int32}

def is_supported_accum(t):
    return t in SUPPORTED_C

When it happens

Trigger: Calling make_instr_desc with c_type such as cutlass.BFloat16 (not in the accepted set), cutlass.Float8E4M3FN, or any integer width other than Int32.

Common situations: Writing an SM100 GEMM with a BF16 accumulator (which UMMA does not encode in this descriptor path); copy-pasting an operand dtype into the accumulator slot when building the descriptor.

Related errors

• Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}
• SplitKV partial output (mO) must be Float32
• All tensors must have the same data type
• Only Float16 or BFloat16 is supported
• LSE tensor must be Float32

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