{"record":{"id":"0d66323df97c11c8","repo":"jax-ml/jax","slug":"tcgen05-mma-only-allows-arriving-on-a-barrier","errorCode":null,"errorMessage":"tcgen05_mma only allows arriving on a Barrier","messagePattern":"tcgen05_mma 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":1646,"sourceCode":"        self.thread,\n        increment_clock=False,\n        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,","sourceCodeStart":1628,"sourceCodeEnd":1664,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L1628-L1664","documentation":"When a tcgen05_mma (tensor-core matmul on Blackwell) op specifies a commit barrier, the interpreter requires it to be a standard memory.Barrier. Any other barrier type is rejected before arriving.","triggerScenarios":"Using a tcgen05 MMA with a barrier_key whose resolved barrier in shared memory is not a memory.Barrier instance (e.g. an mbarrier variant or unsupported type), under interpret mode.","commonSituations":"Experimenting with Blackwell tensor-core pipelines and mixed barrier types; interpreter not covering newer barrier classes.","solutions":["Use a plain Barrier created via standard mosaic APIs for the MMA commit","Ensure the barrier_key references a barrier allocated as memory.Barrier","Update jax for newer tcgen05 interpret support","Run on device if the exotic barrier is required"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"def is_plain_barrier(b) -> bool:\n    from jax._src.pallas.mosaic_gpu.interpret import memory\n    return isinstance(b, memory.Barrier)","tryCatchPattern":"try:\n    kernel(x)\nexcept ValueError as e:\n    if 'tcgen05_mma only allows arriving on a Barrier' in str(e):\n        use standard Barrier for the MMA commit","preventionTips":["Reserve plain Barriers for tcgen05 commit paths in interpret tests"],"tags":["pallas","mosaic-gpu","tcgen05","mma","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"}