{"record":{"id":"8cf01f6348075ccf","repo":"jax-ml/jax","slug":"barrier-offset-offset-is-out-of-bounds","errorCode":null,"errorMessage":"Barrier offset {offset} is out of bounds","messagePattern":"Barrier offset (.+?) is out of bounds","errorType":"validation","errorClass":"IndexError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":1072,"sourceCode":"      nvvm.mbarrier_init(\n          getelementptr(address, [i], i64),\n          c(arrival_count, i32),\n          predicate=predicate,\n      )\n    return BarrierRef(address, c(0, i32), phases, num_barriers)\n\n  def __iter__(self) -> Iterator[\"BarrierRef\"]:\n    if self.num_barriers == 1:\n      yield self\n    else:\n      for offset in range(self.num_barriers):\n        yield self[offset]\n\n  def __getitem__(self, offset: ir.Value | int) -> \"BarrierRef\":\n    i32 = ir.IntegerType.get_signless(32)\n    if isinstance(offset, int):\n      if offset >= self.num_barriers:\n        raise IndexError(f\"Barrier offset {offset} is out of bounds\")\n      offset = c(offset, i32)\n    elif isinstance(offset.type, ir.IndexType):\n      offset = arith.index_castui(i32, offset)\n    elif offset.type != i32:\n      raise ValueError(f\"Expected a dynamic index or an integer, got {offset}\")\n    return BarrierRef(\n        self.base_address,\n        arith.addi(self.offset, offset),\n        self.phases,\n        1,\n    )\n\n  @property\n  def _ptx_scope(self) -> str:\n    if self.base_address.type == ir.Type.parse(\"!llvm.ptr<7>\"):\n      return \"cluster\"\n    return \"cta\"\n","sourceCodeStart":1054,"sourceCodeEnd":1090,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L1054-L1090","documentation":"BarrierGroup.__getitem__ validates static integer offsets against the group's barrier count; offset >= num_barriers raises IndexError. It builds a sub-BarrierRef at a byte offset into the mbarrier array.","triggerScenarios":"group[32] on a group initialized with 32 barriers; indexing with a Python int computed from loop induction that can reach num_barriers.","commonSituations":"Iterating over barriers for multi-stage pipelines with an off-by-one range (range(n+1)); reinitializing with fewer barriers but keeping old hardcoded offsets; confusing barrier index with arrival count.","solutions":["Use range(num_barriers) exactly; assert offset < group.num_barriers in debug","Increase the barrier allocation if the index is legitimately needed (subject to the 32-barrier limit)","Derive indices from the same constant used for allocation"],"exampleFix":"# before\nfor i in range(num_stages + 1):\n    group[i].wait_parity(0)\n# after\nfor i in range(num_stages):\n    group[i].wait_parity(0)","handlingStrategy":"validation","validationCode":"assert 0 <= offset < group.num_barriers, (offset, group.num_barriers)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Iterate range(num_barriers), not range(num_stages + 1)","Derive barrier indices from the same constant as the allocation"],"tags":["jax","mosaic-gpu","barrier","index-out-of-bounds"],"backgroundTag":"index-out-of-bounds","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}