{"record":{"id":"ea47abb94a268306","repo":"jax-ml/jax","slug":"unsupported-accumulator-dtype-acc-dtype","errorCode":null,"errorMessage":"Unsupported accumulator dtype: {acc_dtype}","messagePattern":"Unsupported accumulator dtype: (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/tcgen05.py","lineNumber":75,"sourceCode":"    sparsity_selector: int | None = None,\n) -> ir.Value:\n  f16 = ir.F16Type.get()\n  f32 = ir.F32Type.get()\n  i32 = ir.IntegerType.get_signless(32)\n\n  desc = 0\n  if sparsity_selector is not None:\n    assert 0 <= sparsity_selector < 3\n    desc |= sparsity_selector\n    desc |= 1 << 2  # Enable sparsity\n  if acc_dtype == f16:\n    d_type_val = 0\n  elif acc_dtype == f32:\n    d_type_val = 1\n  elif acc_dtype == i32:\n    d_type_val = 2\n  else:\n    raise NotImplementedError(f\"Unsupported accumulator dtype: {acc_dtype}\")\n  desc |= (d_type_val << 4)  # D type, bits 4-5\n  # Bit 6 is reserved\n  def get_input_encoding(ty):\n    if ty == f16:\n      assert acc_dtype in {f16, f32}\n      return 0\n    elif ty == ir.BF16Type.get():\n      assert acc_dtype == f32\n      return 1\n    elif ty == ir.Float8E4M3FNType.get():\n      assert acc_dtype in {f16, f32}\n      return 0\n    elif ty == ir.Float8E5M2Type.get():\n      assert acc_dtype in {f16, f32}\n      return 1\n    elif ty == ir.IntegerType.get_signless(8):  # Only s8 for now.\n      assert acc_dtype == i32\n      return 1","sourceCodeStart":57,"sourceCodeEnd":93,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/tcgen05.py#L57-L93","documentation":"create_instr_descriptor encodes the MMA accumulator (D) dtype into a 2-bit hardware instruction descriptor field. Only f16 (0), f32 (1) and i32 (2) have encodings; anything else (e.g. f32[i8-packed], bf16) raises NotImplementedError.","triggerScenarios":"Calling mma()/tcgen05 matmul helpers with an accumulator memref of dtype other than f16/f32/i32, e.g. an f32 tensor-core accumulate replaced by bf16.","commonSituations":"Porting Hopper/Warpgroup matmul code to tcgen05 with non-standard accumulator dtypes; using packed f16x2 accumulators before support landed.","solutions":["Use f32 (most common), f16, or i32 accumulators","For integer matmuls, ensure the accumulator is i32","Check the JAX version — newer releases may add encodings"],"exampleFix":"# before\nacc = tmem.alloc(dtype=ir.BFloat16Type.get(), shape=...)\n\n# after\nacc = tmem.alloc(dtype=ir.F32Type.get(), shape=...)","handlingStrategy":"type-guard","validationCode":"supported = {ir.F16Type.get(), ir.F32Type.get(), ir.IntegerType.get_signless(32)}\nassert acc_dtype in supported, f'unsupported accumulator {acc_dtype}'","typeGuard":"def is_supported_acc_dtype(dt) -> bool:\n    return dt in (ir.F16Type.get(), ir.F32Type.get(), ir.IntegerType.get_signless(32))","tryCatchPattern":null,"preventionTips":["Default to f32 accumulators in tcgen05 kernels","Validate accumulator dtype before building descriptors"],"tags":["tcgen05","matmul","mosaic","gpu","dtype"],"backgroundTag":"unsupported-dtype","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}