{"record":{"id":"cd68e65cab568fcd","repo":"jax-ml/jax","slug":"out-of-bounds-swap-of-allocation-key-swapping","errorCode":null,"errorMessage":"Out-of-bounds swap of {allocation_key}: swapping [{read_write_range}] but buffer has shape {shape} .","messagePattern":"Out-of-bounds swap of (.+?): swapping \\[(.+?)\\] but buffer has shape (.+?) \\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py","lineNumber":677,"sourceCode":"    assert mask.shape == val.shape\n\n  shared_memory = _get_shared_memory()\n\n  read_write_range = interpret_utils.to_range(transforms)\n  ret, (shape, _), clock_ = shared_memory.swap_buffer_content(\n      allocation_key,\n      read_write_range,\n      np.array(val),\n      np.array(mask) if mask is not None else None,\n      thread,\n      increment_clock=increment_clock,\n      logging_info=memory.GPULoggingInfo(mesh_location, thread, source_info),\n  )\n  clock = clock if clock is not None else clock_\n\n  if ret is None:\n    if mask is None:\n      raise ValueError(\n          f\"Out-of-bounds swap of {allocation_key}:\"\n          f\" swapping [{read_write_range}] but buffer has shape\"\n          f\" {shape} .\"\n      )\n    else:\n      # TODO(jburnim): Include indices of out-of-bounds locations where mask\n      # is True.\n      raise ValueError(\n          f\"Out-of-bounds masked swap of {allocation_key}: swapping\"\n          f\" [{read_write_range}] but buffer has shape {shape} . \"\n      )\n\n  if shared_memory.detect_races:\n    assert clock is not None\n    get_races().check_write(\n        thread,\n        clock.generic_clock,\n        allocation_key,","sourceCodeStart":659,"sourceCodeEnd":695,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L659-L695","documentation":"The interpreter's swap (atomic read-modify-write) callback detected that the range being swapped falls outside the buffer shape and no mask was provided, so the access cannot be made safe.","triggerScenarios":"Calling a swap/atomic op (e.g. atomic_add on a Block) in interpret mode where the index range exceeds the buffer; specifically the unmasked branch (mask is None).","commonSituations":"Boundary iterations with non-divisible shapes and unmasked atomics; wrong grid size causing out-of-range swaps on accumulator buffers.","solutions":["Add a mask to the swap/atomic operation so boundary lanes are disabled","Fix the grid to ceil-div and ensure the swap index arithmetic stays in bounds","Pad the target buffer to a multiple of the block shape","Reproduce with small shapes in interpret mode to find the offending iteration"],"exampleFix":"# before\ndst[ds[start, bs]].atomic_add(val)  # unmasked\n# after\nmask = (start + jnp.arange(bs)) < n\ndst[ds[start, bs]].atomic_add(val, mask=mask)","handlingStrategy":"validation","validationCode":"assert all(start + bs <= dim for start, dim in zip(starts, dst.shape))","typeGuard":null,"tryCatchPattern":"try:\n    kernel(x)\nexcept ValueError as e:\n    if 'Out-of-bounds swap' in str(e):\n        add mask or fix indices","preventionTips":["Prefer masked atomics","Validate accumulator ranges per iteration"],"tags":["pallas","mosaic-gpu","interpret-mode","out-of-bounds","atomics","swap"],"backgroundTag":"buffer-out-of-bounds-write","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}