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
- Use one of the supported operand types: Float16, BFloat16, Float32 (TF32 path), Float8E4M3FN, or Float8E5M2.
- If you need another dtype, extend to_UMMA_format with the correct MXF8F6F4Format/enum entry once CUTLASS supports it.
- 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
- Whitelist dtypes before building UMMA descriptors.
- Keep a matrix of dtype->format mappings in tests to catch unsupported combos early.
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
- Unsupported CUTLASS scalar type for accumulator: {cutlass_ty
- 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/01b984108205fe2d.
Report an issue: GitHub.