{"record":{"id":"fe04bbe8e367ff20","repo":"sgl-project/sglang","slug":"unsupported-cutlass-scalar-type-for-accumulator","errorCode":null,"errorMessage":"Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}","messagePattern":"Unsupported CUTLASS scalar type for accumulator: (.+?)","errorType":"exception","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py","lineNumber":103,"sourceCode":"    # 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# ---------------------------------------------------------------------------\n# The constructor – accepts only CUTLASS scalar classes\n# ---------------------------------------------------------------------------\n\n\ndef make_instr_desc(\n    a_type,  # CUTLASS scalar class, e.g. cutlass.Int8\n    b_type,\n    c_type,\n    M: int,  # 64, 128 or 256\n    N: int,  # 8 … 256 (multiple of 8)\n    a_major: Major,\n    b_major: Major,\n    a_neg: ScaleIn = ScaleIn.One,","sourceCodeStart":85,"sourceCodeEnd":121,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/mma_sm100_desc.py#L85-L121","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Use cutlass.Float32 for floating accumulation, cutlass.Float16, or cutlass.Int32.","For BF16 compute, accumulate in Float32 and convert — do not request a BF16 C format.","Verify the argument order of make_instr_desc so c_type is actually the accumulator."],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"import cutlass\n\nSUPPORTED_C = {cutlass.Float16, cutlass.Float32, cutlass.Int32}\n\ndef is_supported_accum(t):\n    return t in SUPPORTED_C","tryCatchPattern":null,"preventionTips":["Accumulate in Float32 unless the instruction explicitly supports the target C format.","Unit-test descriptor construction for every dtype you plan to run."],"tags":["cutlass","sm100","accumulator","dtype"],"backgroundTag":"unsupported-dtype-mapping","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}