jax-ml/jax · error · ValueError

tcgen05_commit_arrive only allows arriving on a Barrier

Error message

tcgen05_commit_arrive only allows arriving on a Barrier

What it means

The tcgen05_commit_arrive callback in the interpreter only accepts standard memory.Barrier objects. Arriving at any other barrier type raises this ValueError.

Source

Thrown at jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py:1887

    mesh_location: memory.MeshLocation,
    thread: memory.Thread,
    barrier_key_as_array: jax.Array,
    collective_axis: str | None = None,
    source_info: source_info_util.SourceInfo | None = None,
):
  # TODO(paulbib): Support collective_axis.
  del collective_axis
  barrier_key = HostAllocationKey.from_array(barrier_key_as_array)

  shared_memory = _get_shared_memory()
  if shared_memory.detect_races:
    shared_memory.incr_clock(thread)

  def f(tma_thread_id: int):
    shared_memory = _get_shared_memory()
    barrier = shared_memory.get_barrier(barrier_key)
    if not isinstance(barrier, memory.Barrier):
      raise ValueError(
          "tcgen05_commit_arrive only allows arriving on a Barrier"
      )
    if not barrier.orders_tensor_core:
      raise ValueError(
          "tcgen05_commit_arrive only allows arriving on a Barrier that orders"
          " tensor core"
      )

    clock = None
    if shared_memory.detect_races:
      clock = shared_memory.get_clock(thread)
      completions_clock = shared_memory.get_tcgen05_async_clock(thread)
      assert clock is not None
      if completions_clock is not None:
        clock.update(completions_clock)

    barrier.arrive(
        thread,

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Use a standard Barrier for the tcgen05 commit arrive
2. Update jax / mosaic plugin to matching versions
3. Run on device while using unsupported barrier types
4. Report upstream with a repro

def is_plain_barrier(b) -> bool:
    return isinstance(b, memory.Barrier)

try:
    kernel(x)
except ValueError as e:
    if 'tcgen05_commit_arrive only allows arriving on a Barrier' in str(e):
        use standard Barrier

When it happens

Trigger: Calling tcgen05_commit_arrive where the resolved barrier for the given key is not a memory.Barrier (unsupported variant), during interpret-mode execution.

Common situations: Using newer/experimental barrier kinds with tcgen05 commit paths; version skew between kernel and interpreter support.

Related errors

• tcgen05_mma only allows arriving on a Barrier
• Unsupported barrier type: {type(barrier)}
• tcgen05_mma only allows arriving on a Barrier that orders te
• tcgen05_commit_arrive only allows arriving on a Barrier that
• Expected metadata dtype to be uint2, got: {meta.dtype}

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/87e6586805a9ca18. Report an issue: GitHub.