{"record":{"id":"07adb339ab4c45f9","repo":"jax-ml/jax","slug":"tcgen05-commit-arrive-only-allows-arriving-on-a-ba-07adb3","errorCode":null,"errorMessage":"tcgen05_commit_arrive only allows arriving on a Barrier that orders tensor core","messagePattern":"tcgen05_commit_arrive only allows arriving on a Barrier that orders tensor core","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py","lineNumber":1891,"sourceCode":"    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,\n        clock,\n        memory.GPULoggingInfo(mesh_location, thread, source_info),\n    )\n","sourceCodeStart":1873,"sourceCodeEnd":1909,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L1873-L1909","documentation":"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.","triggerScenarios":"Calling tcgen05_commit_arrive on a Barrier created without tensor-core ordering in interpret mode.","commonSituations":"Sharing a load/TMA barrier with the MMA commit path; forgetting to set the ordering flag when allocating the commit barrier.","solutions":["Create the commit barrier with orders_tensor_core=True","Keep a dedicated barrier for tcgen05 commit, separate from load barriers","Mirror the setup from official mosaic_gpu tcgen05 examples","Update jax if the ordering-flag API differs in your version"],"exampleFix":"# before\ncommit_bar = mgpu.Barrier(threads, 1)\n# after\ncommit_bar = mgpu.Barrier(threads, 1, orders_tensor_core=True)","handlingStrategy":"validation","validationCode":"commit_bar = mgpu.Barrier(threads, 1, orders_tensor_core=True)\nassert commit_bar.orders_tensor_core","typeGuard":"def valid_commit_barrier(b) -> bool:\n    return isinstance(b, memory.Barrier) and b.orders_tensor_core","tryCatchPattern":"try:\n    kernel(x)\nexcept ValueError as e:\n    if 'orders tensor core' in str(e):\n        set orders_tensor_core=True on the commit barrier","preventionTips":["Dedicate a tc-ordered barrier for commits"],"tags":["pallas","mosaic-gpu","tcgen05","barrier","tensor-core"],"backgroundTag":"wrong-synchronization-primitive-configuration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}