{"record":{"id":"e477e146db7f5a75","repo":"xai-org/x-algorithm","slug":"unsupported-cutlass-scalar-type-for-a-b-cutlass","errorCode":null,"errorMessage":"Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}","messagePattern":"Unsupported CUTLASS scalar type for A/B: (.+?)","errorType":"validation","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py","lineNumber":132,"sourceCode":"    MaxShift32 = 3\n\n\ndef to_UMMA_format(cutlass_type) -> int:\n    if cutlass_type is cutlass.Int8:\n        return S8Format.INT8\n    if cutlass_type is cutlass.Uint8:\n        return S8Format.UINT8\n    if cutlass_type is cutlass.Float16:\n        return F16F32Format.F16\n    if cutlass_type is cutlass.BFloat16:\n        return F16F32Format.BF16\n    if cutlass_type is cutlass.TFloat32:\n        return F16F32Format.TF32\n    if cutlass_type is cutlass.Float8E4M3FN:\n        return MXF8F6F4Format.E4M3\n    if cutlass_type is cutlass.Float8E5M2:\n        return MXF8F6F4Format.E5M2\n    raise TypeError(f\"Unsupported CUTLASS scalar type for A/B: {cutlass_type!r}\")\n\n\ndef to_C_format(cutlass_type) -> int:\n    if cutlass_type is cutlass.Float16:\n        return CFormat.F16\n    if cutlass_type is cutlass.Float32:\n        return CFormat.F32\n    if cutlass_type is cutlass.Int32:\n        return CFormat.S32\n    raise TypeError(f\"Unsupported CUTLASS scalar type for accumulator: {cutlass_type!r}\")\n\n\ndef make_instr_desc(\n    a_type,\n    b_type,\n    c_type,\n    M: int,\n    N: int,","sourceCodeStart":114,"sourceCodeEnd":150,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/cutedsl/ranker_fa4/mma_sm100_desc.py#L114-L150","documentation":"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.","triggerScenarios":"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.","commonSituations":"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.","solutions":["Use one of the supported A/B dtypes: Float16, BFloat16, Float32/TF32, Float8E4M3FN, or Float8E5M2","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","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"],"exampleFix":"# before\na_type = cutlass.Int8  # unsupported\n# after\na_type = cutlass.Float8E4M3FN","handlingStrategy":"type-guard","validationCode":"_SUPPORTED_AB = {cutlass.Float16, cutlass.BFloat16, cutlass.TFloat32, cutlass.Float32, cutlass.Float8E4M3FN, cutlass.Float8E5M2}\nassert a_type in _SUPPORTED_AB and b_type in _SUPPORTED_AB, f'unsupported A/B type: {a_type}, {b_type}'","typeGuard":"def is_supported_ab_dtype(t) -> bool:\n    return t in {cutlass.Float16, cutlass.BFloat16, cutlass.TFloat32, cutlass.Float32, cutlass.Float8E4M3FN, cutlass.Float8E5M2}","tryCatchPattern":"try:\n    desc = make_instr_desc(a_type, b_type, c_type, M, N)\nexcept TypeError as e:\n    raise ValueError(f'Unsupported MMA dtype config: {e}') from e","preventionTips":["Whitelist dtypes at the op-construction layer, not deep in descriptor building","Log the derived a_type/b_type when auto-inferring dtypes from tensors"],"tags":["cutlass","tcgen05","mma","dtype","sm100"],"backgroundTag":"unsupported-dtype","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}