jax-ml/jax · error · ValueError

tcgen05_commit_arrive only allows arriving on a Barrier that

Error message

tcgen05_commit_arrive only allows arriving on a Barrier that orders tensor core

What it means

tcgen05_commit_arrive requires its barrier to order tensor-core operations (orders_tensor_core=True). The interpreter validates this before arriving and rejects barriers lacking the ordering.

Source

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

    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,
        clock,
        memory.GPULoggingInfo(mesh_location, thread, source_info),
    )

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Create the commit barrier with orders_tensor_core=True
2. Keep a dedicated barrier for tcgen05 commit, separate from load barriers
3. Mirror the setup from official mosaic_gpu tcgen05 examples
4. Update jax if the ordering-flag API differs in your version

Example fix

# before
commit_bar = mgpu.Barrier(threads, 1)
# after
commit_bar = mgpu.Barrier(threads, 1, orders_tensor_core=True)

commit_bar = mgpu.Barrier(threads, 1, orders_tensor_core=True)
assert commit_bar.orders_tensor_core

def valid_commit_barrier(b) -> bool:
    return isinstance(b, memory.Barrier) and b.orders_tensor_core

try:
    kernel(x)
except ValueError as e:
    if 'orders tensor core' in str(e):
        set orders_tensor_core=True on the commit barrier

When it happens

Trigger: Calling tcgen05_commit_arrive on a Barrier created without tensor-core ordering in interpret mode.

Common situations: Sharing a load/TMA barrier with the MMA commit path; forgetting to set the ordering flag when allocating the commit barrier.

Related errors

• tcgen05_mma only allows arriving on a Barrier that orders te
• tcgen05_mma only allows arriving on a Barrier
• tcgen05_commit_arrive only allows arriving on a Barrier
• Expected metadata dtype to be uint2, got: {meta.dtype}
• Expected metadata to be 3-dimensional (M, K // 4, 2), but it

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