{"record":{"id":"cdd1abb5afc2d66d","repo":"jax-ml/jax","slug":"block-id-must-be-zero-when-allocating-a-buffer-f","errorCode":null,"errorMessage":"`block_id` must be zero when allocating a buffer for all threads","messagePattern":"`block_id` must be zero when allocating a buffer for all threads","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py","lineNumber":275,"sourceCode":"    value: Array of values to initialize the allocated buffer with.\n\n  Returns:\n    `AllocationKey` to refer to the allocated buffer.\n\n  Raises:\n    ValueError: If `thread_id`/`block_id` in `allocation_request` is not zero.\n  \"\"\"\n  allocation_request = HostAllocationRequest.from_array(\n      allocation_request_as_array\n  )\n  del allocation_request_as_array\n\n  if allocation_request.thread_id != 0:\n    raise ValueError(\n        \"`thread_id` must be zero when allocating a buffer for all threads\"\n    )\n  if allocation_request.block_id != 0:\n    raise ValueError(\n        \"`block_id` must be zero when allocating a buffer for all threads\"\n    )\n  assert allocation_request.memory_space_id != memory.get_memory_space_idx(\n      mosaic_gpu_core.MemorySpace.REGS\n  )\n\n  shared_memory = _get_shared_memory()\n\n  key: HostAllocationKey | None = None\n  buffer_id: int | None = None\n  for thread in shared_memory.concurrent_threads(device):\n    buffer_id_for_thread_id = shared_memory.get_next_buffer_id(thread)\n    if not buffer_id:\n      buffer_id = buffer_id_for_thread_id\n    else:\n      # We keep the buffer ids in sync across all threads. This implies, in\n      # particular, that every instance of the assignment to `key` below assigns\n      # an `AllocationKey` object with the same attributes.","sourceCodeStart":257,"sourceCodeEnd":293,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L257-L293","documentation":"GPU interpretation-mode counterpart of the thread_id check: HostAllocationRequest.block_id must be 0 for buffers allocated for all threads. An all-thread collective allocation must be anchored at block 0; any other block indicates invalid allocation bookkeeping in the interpreted kernel.","triggerScenarios":"Interpreting a kernel that issues a for-all-threads allocation from a simulated block other than block 0 (non-zero block_id in HostAllocationRequest).","commonSituations":"Simulating multi-block kernels with collective allocations; custom allocation callbacks constructed manually with wrong block_id; version changes in the interpret backend's block simulation.","solutions":["Anchor collective allocations at block 0 / outside per-block mapped code","If constructing HostAllocationRequest manually, set block_id=0 (and thread_id=0) for all-thread allocations","Minimize the kernel to isolate the offending allocation and report upstream if compiler-emitted"],"exampleFix":"# before\nreq = HostAllocationRequest(..., thread_id=tid, block_id=bid)\n\n# after\nreq = HostAllocationRequest(..., thread_id=0, block_id=0)  # collective alloc","handlingStrategy":"validation","validationCode":"assert allocation_request.block_id == 0, 'collective allocation must come from block 0'","typeGuard":null,"tryCatchPattern":"try:\n    run_interpreted(kernel)\nexcept ValueError as e:\n    if 'block_id' in str(e):\n        anchor_collective_allocation_at_block_zero(kernel)\n        run_interpreted(kernel)\n    else:\n        raise","preventionTips":["Anchor collective allocations at block 0","When building HostAllocationRequest manually, zero thread_id/block_id for all-thread scopes","Report compiler-emitted violations upstream with a minimal reproducer"],"tags":["jax","pallas","mosaic-gpu","interpret-mode","allocation","block-id"],"backgroundTag":"invalid-thread-allocation-scope","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}