{"record":{"id":"878e4ce2533ca627","repo":"jax-ml/jax","slug":"accumulators-are-not-available-on-tpu-info-chip-v","errorCode":null,"errorMessage":"Accumulators are not available on TPU {info.chip_version}","messagePattern":"Accumulators are not available on TPU (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic/core.py","lineNumber":223,"sourceCode":"    )\n    object.__setattr__(self, \"opt_level\", opt_level)\n\n  # Replace is a method, not a field.\n  replace = dataclasses.replace\n\n\ndef check_accumulator_ref(shape: tuple[int, ...], dtype: jnp.dtype, mxu_id: int):\n  from jax._src.pallas.mosaic import tpu_info  # pyrefly: ignore[missing-module-attribute]\n  if len(shape) < 2:\n    raise ValueError(f\"Acc ref must be at least 2D, got shape {shape}\")\n\n  if dtype not in (jnp.float32, jnp.int32):\n    raise ValueError(\n        f\"Acc ref dtype must be float32 or int32, got {dtype}\")\n\n  info = tpu_info.get_tpu_info()\n  if not info.num_accumulators:\n    raise ValueError(\n        f\"Accumulators are not available on TPU {info.chip_version}\"\n    )\n\n  if mxu_id < 0 or mxu_id >= info.num_mxus:\n    raise ValueError(f\"mxu_id must be in [0, {info.num_mxus}), got {mxu_id=}\")\n\n  m, n = math.prod(shape[:-1]), shape[-1]\n  if n != info.mxu_column_size:\n    raise ValueError(\n        f\"The minor dimension size of an accumulator ref must be \"\n        f\"{info.mxu_column_size} but got {n}\"\n    )\n  if m <= 0 or m % info.num_sublanes != 0:\n    raise ValueError(\n        f\"The product of the major dimensions must be a multiple of \"\n        f\"{info.num_sublanes}, but got {m}\"\n    )\n","sourceCodeStart":205,"sourceCodeEnd":241,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/core.py#L205-L241","documentation":"The target TPU chip reports zero available accumulators (tpu_info.num_accumulators == 0), meaning the hardware generation does not expose MXU accumulators to Mosaic kernels. Accumulator refs therefore cannot be validated further and are rejected.","triggerScenarios":"Running a kernel using ACC refs on a TPU chip version whose info reports no accumulators (older/edge TPU generations or certain emulated environments).","commonSituations":"Developing on a different TPU generation than production (e.g., v4 vs v5e/v6e class differences); running on simulators/libtpu runtimes that underreport hardware info.","solutions":["Check tpu_info.get_tpu_info() output for your device and pick a kernel variant without accumulators","Use plain VMEM buffers and regular dot ops instead of ACC refs on this chip","Run on a TPU version that exposes accumulators if the kernel requires them"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":"from jax._src.pallas.mosaic import tpu_info\ninfo = tpu_info.get_tpu_info()\nif not info.num_accumulators:\n    use_vmem_fallback = True","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Probe hardware capabilities before selecting the ACC-based kernel path","Maintain an accumulator-free kernel variant for chips without accumulators"],"tags":["jax","pallas","tpu","hardware","accumulator"],"backgroundTag":"unsupported-hardware-feature","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}