{"record":{"id":"0a627788c7517a79","repo":"xai-org/x-algorithm","slug":"unsupported-tcgen05-mma-op-kind-type-op-name","errorCode":null,"errorMessage":"Unsupported tcgen05 MMA op kind: {type(op).__name__}","messagePattern":"Unsupported tcgen05 MMA op kind: (.+?)","errorType":"exception","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/blackwell_helpers.py","lineNumber":60,"sourceCode":"import xrex.cutedsl.ranker_fa4.mma_sm100_desc as sm100_desc\n\n\ndef _tcgen05_mma_kind(op: cute.nvgpu.tcgen05.mma.MmaOp) -> str:\n    if isinstance(op, tcgen05.mma.MmaF16BF16Op):\n        return \"f16\"\n    if isinstance(op, tcgen05.mma.MmaTF32Op):\n        return \"tf32\"\n    if isinstance(op, tcgen05.mma.MmaI8Op):\n        return \"i8\"\n    if isinstance(op, tcgen05.mma.MmaFP8Op):\n        return \"f8f6f4\"\n    if isinstance(op, tcgen05.mma.MmaMXF8Op):\n        return \"mxf8f6f4\"\n    if isinstance(op, tcgen05.mma.MmaMXF4Op):\n        return \"mxf4\"\n    if isinstance(op, tcgen05.mma.MmaMXF4NVF4Op):\n        return \"mxf4nvf4\"\n    raise TypeError(f\"Unsupported tcgen05 MMA op kind: {type(op).__name__}\")\n\n\n@cute.jit\ndef gemm_w_idx(\n    tiled_mma: cute.TiledMma,\n    acc: cute.Tensor,\n    tCrA: cute.Tensor,\n    tCrB: cute.Tensor,\n    A_idx: Optional[Int32] = None,\n    B_idx: Optional[Int32] = None,\n    zero_init: bool | Boolean = False,\n    swap_AB: bool = False,\n    num_unroll_groups: int = 1,\n) -> None:\n    if const_expr(swap_AB):\n        return gemm_w_idx(\n            tiled_mma, acc, tCrB, tCrA, B_idx, A_idx, zero_init=zero_init, swap_AB=False\n        )","sourceCodeStart":42,"sourceCodeEnd":78,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/blackwell_helpers.py#L42-L78","documentation":"_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.","triggerScenarios":"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).","commonSituations":"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.","solutions":["Use one of the supported MX ops: tcgen05.mma.MmaMXF8Op, MmaMXF4Op, or MmaMXF4NVF4Op","If you need a new op kind, extend _tcgen05_mma_kind with an isinstance branch and corresponding PTX lowering","Pin the cutlass/cuda-parallel version this code was developed against"],"exampleFix":"# before\nop = tcgen05.mma.MmaF16Op(...)\nkind = _tcgen05_mma_kind(op)  # TypeError\n# after\nop = tcgen05.mma.MmaMXF8Op(...)\nkind = _tcgen05_mma_kind(op)","handlingStrategy":"type-guard","validationCode":"from phoenix.xrex.cutedsl.ranker_fa4.blackwell_helpers import _tcgen05_mma_kind\nassert _tcgen05_mma_kind(op) in (\"mxf8f6f4\", \"mxf4\", \"mxf4nvf4\")","typeGuard":"def is_supported_mma_op(op) -> bool:\n    import cutlass.cute.tiled_mma as tcgen05\n    return isinstance(op, (tcgen05.mma.MmaMXF8Op, tcgen05.mma.MmaMXF4Op, tcgen05.mma.MmaMXF4NVF4Op))","tryCatchPattern":"try:\n    kind = _tcgen05_mma_kind(op)\nexcept TypeError:\n    raise NotImplementedError(f\"no PTX lowering for {type(op).__name__}; use an MX op\") from None","preventionTips":["Construct TiledMma only from the supported MmaMX* ops","Pin cutlass/CUDA versions; re-run gemm unit tests after upgrades"],"tags":["cuda","cutedsl","tcgen05","mma","unsupported-operation"],"backgroundTag":"unsupported-operation-type","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}