xai-org/x-algorithm · error · TypeError

Unsupported CUTLASS scalar type for accumulator: {cutlass_ty

Error message

Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}

What it means

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.

Source

Thrown at phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py:142

    if cutlass_type is cutlass.BFloat16:
        return F16F32Format.BF16
    if cutlass_type is cutlass.TFloat32:
        return F16F32Format.TF32
    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:
    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}")


def make_instr_desc(
    a_type,
    b_type,
    c_type,
    M: int,
    N: int,
    a_major: Major,
    b_major: Major,
    a_neg: ScaleIn = ScaleIn.One,
    b_neg: ScaleIn = ScaleIn.One,
    c_sat: Saturate = Saturate.False_,
    is_sparse: bool = False,
    max_shift: MaxShift = MaxShift.NoShift,
) -> int:
    a_fmt = int(to_UMMA_format(a_type))
    b_fmt = int(to_UMMA_format(b_type))

View on GitHub (pinned to 24c60942c5)

Solutions

1. Use Float32 as the accumulator and convert to bf16/fp8 in the epilogue/store path instead
2. Or use Float16 if a 16-bit accumulator is acceptable for the op
3. If you believe the hardware supports the format, extend to_C_format with the correct CFormat enum value upstream

Example fix

# before
desc = make_instr_desc(a, b, cutlass.BFloat16)
# after
desc = make_instr_desc(a, b, cutlass.Float32)  # convert in epilogue

_SUPPORTED_C = {cutlass.Float16, cutlass.Float32, cutlass.Int32}
assert c_type in _SUPPORTED_C, f'unsupported accumulator: {c_type}'

def is_supported_accumulator(t) -> bool:
    return t in {cutlass.Float16, cutlass.Float32, cutlass.Int32}

try:
    desc = make_instr_desc(a_type, b_type, c_type, M, N)
except TypeError:
    c_type = cutlass.Float32  # fall back, convert in epilogue
    desc = make_instr_desc(a_type, b_type, c_type, M, N)

When it happens

Trigger: 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.

Common situations: 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.

Related errors

• Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}
• M must be 64, 128 or 256
• N must be a multiple of 8 in the range 8…256
• Unsupported tcgen05 MMA op kind: {type(op).__name__}
• Unexpected swizzle shift – want S==3 for M==4

AI-assisted analysis of xai-org/x-algorithm@24c60942c5 (2026-08-28). Data as JSON: /api/errors/7bc3cfee855ef1dc. Report an issue: GitHub.