xai-org/x-algorithm · 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 tcgen05 UMMA instruction-descriptor A/B operand format. It supports F16, BF16, F32/TF32, Float8E4M3FN and Float8E5M2; any other type (e.g. fp64, int8, fp6/fp4 without MX formats) raises TypeError. It is reached via make_instr_desc when building an MMA instruction descriptor.

Source

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

    MaxShift32 = 3


def to_UMMA_format(cutlass_type) -> int:
    if cutlass_type is cutlass.Int8:
        return S8Format.INT8
    if cutlass_type is cutlass.Uint8:
        return S8Format.UINT8
    if cutlass_type is cutlass.Float16:
        return F16F32Format.F16
    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,

View on GitHub (pinned to 24c60942c5)

Solutions

1. Use one of the supported A/B dtypes: Float16, BFloat16, Float32/TF32, Float8E4M3FN, or Float8E5M2
2. If you need a new format, extend to_UMMA_format with the matching MXF8F6F4Format enum member (e.g. MXF8F6F4Format.E2M1/E3M2/E2M3) and send a patch
3. Check what dtype the host wrapper derives before constructing the MMA op; often the real bug is an unexpected dtype coming from upstream tensor conversion

Example fix

# before
a_type = cutlass.Int8  # unsupported
# after
a_type = cutlass.Float8E4M3FN

_SUPPORTED_AB = {cutlass.Float16, cutlass.BFloat16, cutlass.TFloat32, cutlass.Float32, cutlass.Float8E4M3FN, cutlass.Float8E5M2}
assert a_type in _SUPPORTED_AB and b_type in _SUPPORTED_AB, f'unsupported A/B type: {a_type}, {b_type}'

def is_supported_ab_dtype(t) -> bool:
    return t in {cutlass.Float16, cutlass.BFloat16, cutlass.TFloat32, cutlass.Float32, cutlass.Float8E4M3FN, cutlass.Float8E5M2}

try:
    desc = make_instr_desc(a_type, b_type, c_type, M, N)
except TypeError as e:
    raise ValueError(f'Unsupported MMA dtype config: {e}') from e

When it happens

Trigger: Calling make_instr_desc / mma_op_to_idesc with an a_type or b_type outside the supported set, such as cutlass.Int8, cutlass.Float64, or an MX FP6/FP4 type not handled in the chain.

Common situations: Extending the kernel to new dtypes (int8 attention, MXF4/MXF6 blocks) without adding a mapping, or passing an unregistered custom NumericType from a wrapper that infers dtype dynamically.

Related errors

• Unsupported CUTLASS scalar type for accumulator: {cutlass_ty
• 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__}
• The layout of mdV is wrong

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