{"record":{"id":"87e6586805a9ca18","repo":"jax-ml/jax","slug":"tcgen05-commit-arrive-only-allows-arriving-on-a-ba","errorCode":null,"errorMessage":"tcgen05_commit_arrive only allows arriving on a Barrier","messagePattern":"tcgen05_commit_arrive only allows arriving on a Barrier","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py","lineNumber":1887,"sourceCode":"    mesh_location: memory.MeshLocation,\n    thread: memory.Thread,\n    barrier_key_as_array: jax.Array,\n    collective_axis: str | None = None,\n    source_info: source_info_util.SourceInfo | None = None,\n):\n  # TODO(paulbib): Support collective_axis.\n  del collective_axis\n  barrier_key = HostAllocationKey.from_array(barrier_key_as_array)\n\n  shared_memory = _get_shared_memory()\n  if shared_memory.detect_races:\n    shared_memory.incr_clock(thread)\n\n  def f(tma_thread_id: int):\n    shared_memory = _get_shared_memory()\n    barrier = shared_memory.get_barrier(barrier_key)\n    if not isinstance(barrier, memory.Barrier):\n      raise ValueError(\n          \"tcgen05_commit_arrive only allows arriving on a Barrier\"\n      )\n    if not barrier.orders_tensor_core:\n      raise ValueError(\n          \"tcgen05_commit_arrive only allows arriving on a Barrier that orders\"\n          \" tensor core\"\n      )\n\n    clock = None\n    if shared_memory.detect_races:\n      clock = shared_memory.get_clock(thread)\n      completions_clock = shared_memory.get_tcgen05_async_clock(thread)\n      assert clock is not None\n      if completions_clock is not None:\n        clock.update(completions_clock)\n\n    barrier.arrive(\n        thread,","sourceCodeStart":1869,"sourceCodeEnd":1905,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L1869-L1905","documentation":"The tcgen05_commit_arrive callback in the interpreter only accepts standard memory.Barrier objects. Arriving at any other barrier type raises this ValueError.","triggerScenarios":"Calling tcgen05_commit_arrive where the resolved barrier for the given key is not a memory.Barrier (unsupported variant), during interpret-mode execution.","commonSituations":"Using newer/experimental barrier kinds with tcgen05 commit paths; version skew between kernel and interpreter support.","solutions":["Use a standard Barrier for the tcgen05 commit arrive","Update jax / mosaic plugin to matching versions","Run on device while using unsupported barrier types","Report upstream with a repro"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"def is_plain_barrier(b) -> bool:\n    return isinstance(b, memory.Barrier)","tryCatchPattern":"try:\n    kernel(x)\nexcept ValueError as e:\n    if 'tcgen05_commit_arrive only allows arriving on a Barrier' in str(e):\n        use standard Barrier","preventionTips":["Use standard Barriers for tcgen05 commit in interpret mode"],"tags":["pallas","mosaic-gpu","tcgen05","barrier","interpret-mode"],"backgroundTag":"unsupported-operation-in-interpreter","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}