xai-org/x-algorithm · error · TypeError

Unsupported tcgen05 MMA op kind: {type(op).__name__}

Error message

Unsupported tcgen05 MMA op kind: {type(op).__name__}

What it means

_tcgen05_mma_kind maps a tcgen05 MMA op object to a string kind ('mxf8f6f4', 'mxf4', 'mxf4nvf4') used by the PTX-emitting gemm helpers. If the op is not one of the recognized MmaMX* classes it raises TypeError, since the codegen has no lowering for that op type. This is a programming/library-extension error rather than a data error.

Source

Thrown at phoenix/xrex/cutedsl/ranker_fa4/blackwell_helpers.py:60

import xrex.cutedsl.ranker_fa4.mma_sm100_desc as sm100_desc


def _tcgen05_mma_kind(op: cute.nvgpu.tcgen05.mma.MmaOp) -> str:
    if isinstance(op, tcgen05.mma.MmaF16BF16Op):
        return "f16"
    if isinstance(op, tcgen05.mma.MmaTF32Op):
        return "tf32"
    if isinstance(op, tcgen05.mma.MmaI8Op):
        return "i8"
    if isinstance(op, tcgen05.mma.MmaFP8Op):
        return "f8f6f4"
    if isinstance(op, tcgen05.mma.MmaMXF8Op):
        return "mxf8f6f4"
    if isinstance(op, tcgen05.mma.MmaMXF4Op):
        return "mxf4"
    if isinstance(op, tcgen05.mma.MmaMXF4NVF4Op):
        return "mxf4nvf4"
    raise TypeError(f"Unsupported tcgen05 MMA op kind: {type(op).__name__}")


@cute.jit
def gemm_w_idx(
    tiled_mma: cute.TiledMma,
    acc: cute.Tensor,
    tCrA: cute.Tensor,
    tCrB: cute.Tensor,
    A_idx: Optional[Int32] = None,
    B_idx: Optional[Int32] = None,
    zero_init: bool | Boolean = False,
    swap_AB: bool = False,
    num_unroll_groups: int = 1,
) -> None:
    if const_expr(swap_AB):
        return gemm_w_idx(
            tiled_mma, acc, tCrB, tCrA, B_idx, A_idx, zero_init=zero_init, swap_AB=False
        )

View on GitHub (pinned to 24c60942c5)

Solutions

1. Use one of the supported MX ops: tcgen05.mma.MmaMXF8Op, MmaMXF4Op, or MmaMXF4NVF4Op
2. If you need a new op kind, extend _tcgen05_mma_kind with an isinstance branch and corresponding PTX lowering
3. Pin the cutlass/cuda-parallel version this code was developed against

Example fix

# before
op = tcgen05.mma.MmaF16Op(...)
kind = _tcgen05_mma_kind(op)  # TypeError
# after
op = tcgen05.mma.MmaMXF8Op(...)
kind = _tcgen05_mma_kind(op)

from phoenix.xrex.cutedsl.ranker_fa4.blackwell_helpers import _tcgen05_mma_kind
assert _tcgen05_mma_kind(op) in ("mxf8f6f4", "mxf4", "mxf4nvf4")

def is_supported_mma_op(op) -> bool:
    import cutlass.cute.tiled_mma as tcgen05
    return isinstance(op, (tcgen05.mma.MmaMXF8Op, tcgen05.mma.MmaMXF4Op, tcgen05.mma.MmaMXF4NVF4Op))

try:
    kind = _tcgen05_mma_kind(op)
except TypeError:
    raise NotImplementedError(f"no PTX lowering for {type(op).__name__}; use an MX op") from None

When it happens

Trigger: Constructing gemm_ptx / gemm_ptx_loop / gemm_ptx_partial / mma with a TiledMma built from an unsupported tcgen05 op (e.g. a plain MmaF16Op, MmaBF16Op, or a new op class added in a newer cutlass/cuda version).

Common situations: Upgrading NVIDIA cutlass/CUDA extensions that introduce new tcgen05 op kinds not yet handled here; writing a new gemm path with fp16/bf16 ops where only MX fp8/fp4 ops are supported.

Related errors

• block-sparse arrays cover {bs_num_blocks} m-tiles but the ke
• valid_block_upper and valid_block_lower must be provided tog
• block-sparse arrays cover {bs_num_blocks} m-tiles but the ke
• The layout of mdV is wrong
• Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}

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