{"record":{"id":"01b984108205fe2d","repo":"sgl-project/sglang","slug":"unsupported-cutlass-scalar-type-for-a-b-cutlass","errorCode":null,"errorMessage":"Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}","messagePattern":"Unsupported CUTLASS scalar type for A/B: (.+?)","errorType":"exception","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":90,"sourceCode":"    if cutlass_type is cutlass.Int8:\n        return S8Format.INT8\n    # Unsigned 8-bit (if available in your CUTLASS build)\n    if cutlass_type is cutlass.Uint8:\n        return S8Format.UINT8\n    # FP-16 / BF-16\n    if cutlass_type is cutlass.Float16:\n        return F16F32Format.F16\n    if cutlass_type is cutlass.BFloat16:\n        return F16F32Format.BF16\n    # TensorFloat-32 (8-bit exponent, 10-bit mantissa packed in 19 bits)\n    if cutlass_type is cutlass.TFloat32:\n        return F16F32Format.TF32\n    # Float-8 / Float-6 / Float-4 – add whenever CUTLASS exposes them\n    if cutlass_type is cutlass.Float8E4M3FN:\n        return MXF8F6F4Format.E4M3\n    if cutlass_type is cutlass.Float8E5M2:\n        return MXF8F6F4Format.E5M2\n    raise TypeError(f\"Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}\")\n\n\ndef to_C_format(cutlass_type) -> int:\n    \"\"\"\n    Map a CUTLASS scalar class to the 2-bit accumulator encoding.\n    \"\"\"\n    if cutlass_type is cutlass.Float16:\n        return CFormat.F16\n    if cutlass_type is cutlass.Float32:\n        return CFormat.F32\n    if cutlass_type is cutlass.Int32:\n        return CFormat.S32\n    raise TypeError(\n        f\"Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}\"\n    )\n\n\n# ---------------------------------------------------------------------------","sourceCodeStart":72,"sourceCodeEnd":108,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L72-L108","documentation":"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.","triggerScenarios":"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).","commonSituations":"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.","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."],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"import cutlass\n\nSUPPORTED_AB = {\n    cutlass.Float16, cutlass.BFloat16, cutlass.Float32,\n    cutlass.Float8E4M3FN, cutlass.Float8E5M2,\n}\n\ndef is_supported_ab(t):\n    return t in SUPPORTED_AB","tryCatchPattern":null,"preventionTips":["Whitelist dtypes before building UMMA descriptors.","Keep a matrix of dtype->format mappings in tests to catch unsupported combos early."],"tags":["cutlass","sm100","dtype","tensor-core"],"backgroundTag":"unsupported-dtype-mapping","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}