{"record":{"id":"3f258f2ba07fb938","repo":"jax-ml/jax","slug":"mxu-id-must-be-in-0-info-num-mxus-got-mxu-i","errorCode":null,"errorMessage":"mxu_id must be in [0, {info.num_mxus}), got {mxu_id=}","messagePattern":"mxu_id must be in \\[0, (.+?)\\), got (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic/core.py","lineNumber":228,"sourceCode":"\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\n\nclass MemoryRef(pallas_core.MemoryRef):\n\n  def __matmul__(self, other, /):\n    if not isinstance(other, pallas_core.Mesh):","sourceCodeStart":210,"sourceCodeEnd":246,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/core.py#L210-L246","documentation":"Accumulator refs are per-MXU resources; mxu_id must index a valid MXU on the chip (0 <= mxu_id < info.num_mxus). Values outside that range (including negatives) fail validation.","triggerScenarios":"Passing an out-of-range mxu_id when creating an accumulator ref, e.g., mxu_id=4 on a chip with 2 MXUs, or defaulting to -1 as 'unspecified'.","commonSituations":"Hardcoding mxu_id values tuned for one TPU generation and running on another; looping over core indices with off-by-one errors.","solutions":["Query info.num_mxus via tpu_info.get_tpu_info() and clamp/validate mxu_id against it","Use device-derived indices rather than hardcoded constants","Fix off-by-one loops generating mxu_id values"],"exampleFix":"// before\nacc = make_acc(..., mxu_id=8)\n// after\ninfo = tpu_info.get_tpu_info()\nassert 0 <= mxu_id < info.num_mxus\nacc = make_acc(..., mxu_id=mxu_id % info.num_mxus)","handlingStrategy":"validation","validationCode":"from jax._src.pallas.mosaic import tpu_info\ninfo = tpu_info.get_tpu_info()\nassert 0 <= mxu_id < info.num_mxus, (mxu_id, info.num_mxus)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Never hardcode mxu_id; derive from device info","Validate loop bounds that generate mxu_id values"],"tags":["jax","pallas","tpu","accumulator","index-out-of-range"],"backgroundTag":"index-out-of-range","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}