{"record":{"id":"4484caabbd478ff2","repo":"sgl-project/sglang","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":"python/sglang/kernels/ops/attention/flash_attn/cute/blackwell_helpers.py","lineNumber":30,"sourceCode":"\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    # cutlass-dsl >=4.5.2 builds plain FP8 MMAs as MmaF8F6F4Op (make_trivial_tiled_mma's\n    # _F8F6F4_TYPES branch); <4.4.x returned the now-legacy MmaFP8Op. Both map to kind::f8f6f4.\n    if isinstance(op, (tcgen05.mma.MmaFP8Op, tcgen05.mma.MmaF8F6F4Op)):\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":12,"sourceCodeEnd":48,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/flash_attn/cute/blackwell_helpers.py#L12-L48","documentation":"_tcgen05_mma_kind maps a CUTLASS tcgen05 MMA op object to a string kind ('mxf8f6f4', 'mxf4', 'mxf4nvf4') used by the PTX gemm helpers. Any other op type (e.g. a plain fp16/bf16 MmaOp) is unsupported and raises TypeError.","triggerScenarios":"Passing a cute.TiledMma built with an MMA op class not in {MmaMXF8Op, MmaMXF4Op, MmaMXF4NVF4Op} to gemm_ptx/gemm_ptx_loop/gemm_ptx_partial/mma in the Blackwell CUTE helpers.","commonSituations":"Extending FA4 CUTE kernels on Blackwell with non-MX (bf16/fp16/fp8-legacy) MMA configurations; constructing TiledMma with default ops.","solutions":["Build the TiledMma with one of the supported MX op classes (tcgen05.mma.MmaMXF8Op / MmaMXF4Op / MmaMXF4NVF4Op)","Use the existing non-PTX gemm path for unsupported dtypes","Add a new kind mapping in _tcgen05_mma_kind if you control the kernel tree"],"exampleFix":"# before\nop = tcgen05.mma.MmaF16Op()  # unsupported\ntiled_mma = cute.make_tiled_mma(op)\n# after\nop = tcgen05.mma.MmaMXF8Op(...)\ntiled_mma = cute.make_tiled_mma(op)","handlingStrategy":"type-guard","validationCode":"from flash_attn.cute.blackwell_helpers import _tcgen05_mma_kind\\n_tcgen05_mma_kind(tiled_mma.op)  # raises early if unsupported","typeGuard":"def is_supported_mma_op(op) -> bool:\\n    import cutlass.cute.tcgen05.mma as m\\n    return isinstance(op, (m.MmaMXF8Op, m.MmaMXF4Op, m.MmaMXF4NVF4Op))","tryCatchPattern":null,"preventionTips":["Construct TiledMma only from the supported MX op classes","Validate op type before entering gemm_ptx helpers"],"tags":["cutlass","cute","tcgen05","blackwell","mma","unsupported-dtype"],"backgroundTag":"unsupported-mma-op-kind","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}