{"record":{"id":"1de6ede200598e08","repo":"jax-ml/jax","slug":"tcgen05-mma-only-allows-arriving-on-a-barrier-that","errorCode":null,"errorMessage":"tcgen05_mma only allows arriving on a Barrier that orders tensor core","messagePattern":"tcgen05_mma 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":1648,"sourceCode":"        logging_info=logging_info,\n    )\n\n    if shared_memory.detect_races:\n      assert clock is not None\n      get_races().check_write(\n          self.thread,\n          clock.generic_clock,\n          self.acc_key,\n          acc_range,\n          source_info=self.source_info,\n      )\n\n    if self.barrier_key:\n      barrier = shared_memory.get_barrier(self.barrier_key)\n      if not isinstance(barrier, memory.Barrier):\n        raise ValueError(\"tcgen05_mma only allows arriving on a Barrier\")\n      if not barrier.orders_tensor_core:\n        raise ValueError(\n            \"tcgen05_mma only allows arriving on a Barrier that orders tensor\"\n            \" core\"\n        )\n      barrier.arrive(\n          thread=self.thread,\n          clock=clock,\n          logging_info=logging_info,\n      )\n\n    return clock.copy() if clock is not None else None\n\n\ndef tcgen05_mma(\n    *,\n    token: jax.Array,\n    mesh_location: memory.MeshLocation,\n    thread: memory.Thread,\n    acc_allocation_key_as_array: jax.Array,","sourceCodeStart":1630,"sourceCodeEnd":1666,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L1630-L1666","documentation":"The interpreter requires that the Barrier a tcgen05_mma arrives on has orders_tensor_core=True, i.e. it must order tensor-core (async MMA) operations, otherwise synchronization semantics would be wrong.","triggerScenarios":"Attaching a commit barrier to tcgen05_mma where the Barrier was created without tensor-core ordering (orders_tensor_core flag False), in interpret mode.","commonSituations":"Reusing a generic barrier intended for TMA/loads as an MMA commit barrier; missing the ordering option when allocating the barrier.","solutions":["Create the barrier with tensor-core ordering enabled (e.g. Barrier(..., orders_tensor_core=True) or the mosaic API that sets it)","Allocate a dedicated commit barrier for MMA separate from load barriers","Update jax — API names for ordering flags have changed across versions","Check mosaic_gpu examples for the canonical tcgen05 pipeline barrier setup"],"exampleFix":"# before\nbarrier = mgpu.Barrier(threads, num_workers=1)  # no tc ordering\n# after\nbarrier = mgpu.Barrier(threads, num_workers=1, orders_tensor_core=True)","handlingStrategy":"validation","validationCode":"bar = mgpu.Barrier(threads, 1, orders_tensor_core=True)\nassert bar.orders_tensor_core","typeGuard":"def orders_tc(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        recreate barrier with orders_tensor_core=True","preventionTips":["Always set orders_tensor_core on MMA commit barriers"],"tags":["pallas","mosaic-gpu","tcgen05","mma","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"}