{"record":{"id":"c5912c5345778d2b","repo":"jax-ml/jax","slug":"barriers-are-not-arrays","errorCode":null,"errorMessage":"Barriers are not arrays","messagePattern":"Barriers are not arrays","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":1511,"sourceCode":"      guarantee that the TensorCore-related operations in other threads have\n      completed. Similarly, when False any TensorCore operation in the waiting\n      thread is allowed to begin before the wait succeeds.\n  \"\"\"\n  num_arrivals: int = 1\n  num_barriers: int | Sequence[int] = 1\n  orders_tensor_core: bool = False\n\n  def __post_init__(self):\n    if (n := self.num_arrivals) < 1:\n      raise ValueError(f\"Num arrivals must be at least 1, but got {n}\")\n\n    if isinstance(self.num_barriers, int):\n      object.__setattr__(self, \"num_barriers\", (self.num_barriers,))\n    else:\n      object.__setattr__(self, \"num_barriers\", tuple(self.num_barriers))\n\n  def get_array_aval(self) -> jax_core.ShapedArray:\n    raise ValueError(\"Barriers are not arrays\")\n\n  def get_ref_aval(self) -> state.AbstractRef:\n    ty = BarrierType(self.num_arrivals, self.orders_tensor_core)\n    return state.AbstractRef(jax_core.ShapedArray(self.num_barriers, ty), SMEM)\n\n\n@dataclasses.dataclass(frozen=True, kw_only=True)\nclass ClusterBarrier:\n  collective_axes: tuple[str | tuple[str, ...], ...]\n  num_barriers: int | Sequence[int] = 1\n  num_arrivals: int = 1\n  orders_tensor_core: bool = False\n  leader_tracked: bool = False\n\n  def __post_init__(self):\n    if (n := self.num_arrivals) < 1:\n      raise ValueError(f\"Num arrivals must be at least 1, but got {n}\")\n","sourceCodeStart":1493,"sourceCodeEnd":1529,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L1493-L1529","documentation":"BarrierSpec.get_array_aval always raises because a BarrierSpec describes hardware synchronization barriers, not a data array. Barriers have no array semantics; only get_ref_aval (an AbstractRef in SMEM) is meaningful. Calling the array path signals a programming/API misuse.","triggerScenarios":"Code that treats a BarrierSpec as an output/input array — e.g. passing a barrier spec where an out_shape/array aval is expected, or generic code that calls get_array_aval() on every spec.","commonSituations":"Generic plumbing in user kernels or JAX transforms that iterates specs and requests array avals unconditionally; confusing barrier buffers with regular SMEM arrays.","solutions":["Branch on the spec type: use get_ref_aval() for BarrierSpec/ClusterBarrierSpec, get_array_aval() only for array specs","Don't pass BarrierSpec where an out_shape/array is required","Keep barriers out of any code path that computes array shapes/dtypes"],"exampleFix":"// before\naval = spec.get_array_aval()  # raises for barriers\n// after\naval = spec.get_ref_aval() if isinstance(spec, plgpu.BarrierSpec) else spec.get_array_aval()","handlingStrategy":"type-guard","validationCode":"if isinstance(spec, plgpu.BarrierSpec):\n    aval = spec.get_ref_aval()\nelse:\n    aval = spec.get_array_aval()","typeGuard":"def is_barrier_spec(spec) -> bool:\\n    return isinstance(spec, plgpu.BarrierSpec)","tryCatchPattern":null,"preventionTips":["Dispatch on spec type before requesting avals","Never pass barrier specs as out_shape"],"tags":["jax","pallas","barrier","api-misuse"],"backgroundTag":"wrong-aval-kind-requested","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}