sgl-project/sglang · error · TypeError

Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}

Error message

Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}

What it means

to_UMMA_format maps a CUTLASS scalar type to the UMMA instruction descriptor's A/B operand format. Only F16, BF16, TF32, Float8E4M3FN and Float8E5M2 are supported; any other scalar class passed to make_instr_desc raises this TypeError.

Source

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

    if cutlass_type is cutlass.Int8:
        return S8Format.INT8
    # Unsigned 8-bit (if available in your CUTLASS build)
    if cutlass_type is cutlass.Uint8:
        return S8Format.UINT8
    # FP-16 / BF-16
    if cutlass_type is cutlass.Float16:
        return F16F32Format.F16
    if cutlass_type is cutlass.BFloat16:
        return F16F32Format.BF16
    # TensorFloat-32 (8-bit exponent, 10-bit mantissa packed in 19 bits)
    if cutlass_type is cutlass.TFloat32:
        return F16F32Format.TF32
    # 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}"
    )


# ---------------------------------------------------------------------------

View on GitHub (pinned to 0132848349)

Solutions

  1. Use one of the supported operand types: Float16, BFloat16, Float32 (TF32 path), Float8E4M3FN, or Float8E5M2.
  2. If you need another dtype, extend to_UMMA_format with the correct MXF8F6F4Format/enum entry once CUTLASS supports it.
  3. Check that you are not accidentally passing c_type (accumulator) where a_type is expected.
Defensive patterns

Strategy: type-guard

Type guard

import cutlass

SUPPORTED_AB = {
    cutlass.Float16, cutlass.BFloat16, cutlass.Float32,
    cutlass.Float8E4M3FN, cutlass.Float8E5M2,
}

def is_supported_ab(t):
    return t in SUPPORTED_AB

Prevention

When it happens

Trigger: Constructing an SM100 UMMA instruction descriptor (make_instr_desc) with a_type/b_type like cutlass.Float8E8M0FNU, cutlass.Int8, cutlass.Float64, or an FP6/FP4 type (explicitly noted as unsupported until CUTLASS exposes them).

Common situations: Extending the SM100 GEMM path to new dtypes (MXFP4/FP6, INT8) without adding the mapping; passing the accumulator type by mistake as an operand type.

Related errors


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