{"record":{"id":"7bc3cfee855ef1dc","repo":"xai-org/x-algorithm","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":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":142,"sourceCode":"    if cutlass_type is cutlass.BFloat16:\n        return F16F32Format.BF16\n    if cutlass_type is cutlass.TFloat32:\n        return F16F32Format.TF32\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    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(f\"Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}\")\n\n\ndef make_instr_desc(\n    a_type,\n    b_type,\n    c_type,\n    M: int,\n    N: int,\n    a_major: Major,\n    b_major: Major,\n    a_neg: ScaleIn = ScaleIn.One,\n    b_neg: ScaleIn = ScaleIn.One,\n    c_sat: Saturate = Saturate.False_,\n    is_sparse: bool = False,\n    max_shift: MaxShift = MaxShift.NoShift,\n) -> int:\n    a_fmt = int(to_UMMA_format(a_type))\n    b_fmt = int(to_UMMA_format(b_type))","sourceCodeStart":124,"sourceCodeEnd":160,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L124-L160","documentation":"to_C_format maps the CUTLASS accumulator (C/D) scalar type to the UMMA descriptor C format, supporting only Float16, Float32, and Int32. Any other accumulator type raises TypeError. It is called from make_instr_desc when building an instruction descriptor.","triggerScenarios":"Calling make_instr_desc with a c_type such as cutlass.BFloat16 or cutlass.Float8E4M3FN as the accumulator, which has no hardware C-format encoding on SM100.","commonSituations":"Trying to fuse the epilogue so the MMA writes bf16/fp8 output directly, or copying an MMA op definition from another architecture (SM90) whose accumulator type differs.","solutions":["Use Float32 as the accumulator and convert to bf16/fp8 in the epilogue/store path instead","Or use Float16 if a 16-bit accumulator is acceptable for the op","If you believe the hardware supports the format, extend to_C_format with the correct CFormat enum value upstream"],"exampleFix":"# before\ndesc = make_instr_desc(a, b, cutlass.BFloat16)\n# after\ndesc = make_instr_desc(a, b, cutlass.Float32)  # convert in epilogue","handlingStrategy":"type-guard","validationCode":"_SUPPORTED_C = {cutlass.Float16, cutlass.Float32, cutlass.Int32}\nassert c_type in _SUPPORTED_C, f'unsupported accumulator: {c_type}'","typeGuard":"def is_supported_accumulator(t) -> bool:\n    return t in {cutlass.Float16, cutlass.Float32, cutlass.Int32}","tryCatchPattern":"try:\n    desc = make_instr_desc(a_type, b_type, c_type, M, N)\nexcept TypeError:\n    c_type = cutlass.Float32  # fall back, convert in epilogue\n    desc = make_instr_desc(a_type, b_type, c_type, M, N)","preventionTips":["Default accumulators to Float32 and do output conversion in the epilogue","Validate the accumulator type where the MMA op object is created, closer to user code"],"tags":["cutlass","tcgen05","mma","accumulator","dtype"],"backgroundTag":"unsupported-dtype","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}