{"record":{"id":"de6c7068595f7741","repo":"jax-ml/jax","slug":"mma-with-element-type-elem-type-str-does-not-sup","errorCode":null,"errorMessage":"MMA with element type {elem_type_str} does not support block scaling","messagePattern":"MMA with element type (.+?) does not support block scaling","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":294,"sourceCode":"        raise ValueError(\n            f\"Swizzle={b_swizzle} is too big for MMA with M=64. Try\"\n            \" lowering it.\"\n        )\n  else:\n    raise ValueError(f\"Only M=128 and M=64 are supported for MMA, but got M={m}\")\n  f32 = ir.F32Type.get()\n  f16 = ir.F16Type.get()\n  s32 = ir.IntegerType.get_signless(32)\n  elem_type_str = (\n      f\"{a_element_type}\"\n      if a_element_type == b_element_type\n      else f\"({a_element_type}, {b_element_type})\"\n  )\n  if a_element_type == f32 or a_element_type == ir.BF16Type.get():\n    if a_element_type == f32 and is_sparse:\n      raise NotImplementedError(\"Sparse MMA unsupported for f32\")\n    if is_scaled:\n      raise ValueError(\n          f\"MMA with element type {elem_type_str} does not support block scaling\"\n      )\n    if d.dtype != f32:\n      raise ValueError(\n          f\"MMA with element type {elem_type_str} only supports accumulators\"\n          f\" of type f32, but got: {d.dtype}\"\n      )\n  elif a_element_type == f16:\n    if is_scaled:\n      raise ValueError(\n          f\"MMA with element type {elem_type_str} does not support block scaling\"\n      )\n    if d.dtype != f16 and d.dtype != f32:\n      raise ValueError(\n          f\"MMA with element type {elem_type_str} only supports accumulators of\"\n          f\" type f32 or f16, but got: {d.dtype}\"\n      )\n  elif any(","sourceCodeStart":276,"sourceCodeEnd":312,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L276-L312","documentation":"Block-scaled MMA ( tcgen05 with scale factors, used for MX formats) is only supported for fp8 (and f16 in some modes) operand types. Requesting is_scaled with f32 or bf16 A operands has no hardware instruction, so it is rejected.","triggerScenarios":"Calling tcgen05.mma with a_element_type f32 or bf16 while passing scale operands / is_scaled=True.","commonSituations":"Trying to apply microscaling (MXFP) formats to bf16 weights; enabling scaling flags copied from an fp8 MX kernel onto a bf16 GEMM; experimenting with scaled accumulation precision on higher-precision inputs.","solutions":["Switch A/B operands to fp8 (Float8E5M2 or Float8E4M3FN) to use block scaling","If you need bf16/f32 inputs, drop the scale operands and run unscaled MMA","For dynamic-range issues in bf16, consider tmem-based normalization outside the MMA instead of block scaling"],"exampleFix":"# before\ntcgen05.mma(a_bf16, b_bf16, d, scales=(sa, sb))  # raises\n# after\na8, b8 = a_bf16.astype(mxfp8), b_bf16.astype(mxfp8)\ntcgen05.mma(a8, b8, d, scales=(sa, sb))","handlingStrategy":"validation","validationCode":"if is_scaled:\n    ok = isinstance(a_element_type, (ir.Float8E5M2Type, ir.Float8E4M3FNType)) or a_element_type == ir.F16Type.get()\n    assert ok, f'block scaling unsupported for {a_element_type}'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Reserve block scaling for fp8 MX kernels","Assert scale operand presence matches operand dtype before the call"],"tags":["jax","mosaic","gpu","block-scaling","dtype","mma","tcgen05","mxfp"],"backgroundTag":"unsupported-dtype-for-operation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}