{"record":{"id":"8d8dc8b5cb8747ba","repo":"jax-ml/jax","slug":"mma-barrier-must-have-orders-tensor-core-set-to-tr","errorCode":null,"errorMessage":"MMA barrier must have orders_tensor_core set to True.","messagePattern":"MMA barrier must have orders_tensor_core set to True\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/primitives.py","lineNumber":2591,"sourceCode":"\n  if collective_axis is not None:\n    # TODO(justinfu): If under a core_map, the avals for acc/a\n    # become normal MemRefs so we cannot check if they are collective.\n    # Figure out a way to fix this.\n    if isinstance(acc, gpu_core.AbstractTMEMRef) and not acc.collective:\n      raise ValueError(\n          \"Accumulator Ref must be collective if collective_axis is set.\")\n    if isinstance(a, gpu_core.AbstractTMEMRef) and not a.collective:\n      raise ValueError(\n          \"LHS Ref must be collective if collective_axis is set.\")\n\n  scales_and_transforms_leaves = barrier_scales_and_transforms_leaves\n  if arrive:\n    barrier, *scales_and_transforms_leaves = barrier_scales_and_transforms_leaves\n    orders_tensor_core = getattr(\n        barrier.inner_aval.dtype, \"orders_tensor_core\", False)\n    if not orders_tensor_core:\n      raise ValueError(\"MMA barrier must have orders_tensor_core set to True.\")\n  if scaled:\n    a_scale, b_scale = scales_and_transforms_leaves[:2]\n    if a_scale.memory_space != gpu_core.TMEM:\n      raise ValueError(\"a_scale must be a TMEM Ref\")\n    if b_scale.memory_space != gpu_core.TMEM:\n      raise ValueError(\"b_scale must be a TMEM Ref\")\n\n  return [], {gpu_core._memory_effect}\n\n\n@lowering.register_lowering_rule(tcgen05_mma_p, *gpu_core.LANExWG_SEMANTICS)\n@lowering.register_lowering_rule(tcgen05_mma_p, *gpu_core.LANExWARP_SEMANTICS)\ndef _tcgen05_mma_lowering(\n    ctx: lowering.LoweringRuleContext,\n    acc: tcgen05.TMEMRef,\n    a_ref,\n    b_ref,\n    accumulate: bool | ir.Value,","sourceCodeStart":2573,"sourceCodeEnd":2609,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/primitives.py#L2573-L2609","documentation":"When arrive=True in tcgen05.mma, the provided MMA barrier must have orders_tensor_core=True on its dtype so the barrier arrival is ordered by the tensor core, not the SM. This is read via getattr on barrier.inner_aval.dtype.","triggerScenarios":"Passing a regular SemaphoreBarrier (orders_tensor_core defaults False) as the barrier with arrive=True, e.g. tcgen05.mma(..., barrier=bar, arrive=True).","commonSituations":"Reusing a barrier created for loads/stores in an MMA pipeline; forgetting to construct the barrier with tensor-core ordering when doing software pipelining with tcgen05.","solutions":["Create the barrier with orders_tensor_core=True (mosaic_gpu barrier allocation option)","Or drop arrive=True and manage the arrival manually","Check getattr(bar.dtype, 'orders_tensor_core', False) before the call"],"exampleFix":"# before\nbar = allocate_barrier()\ntcgen05.mma(a, b, acc, k_dim=k, barrier=bar, arrive=True)\n# after\nbar = allocate_barrier(orders_tensor_core=True)\ntcgen05.mma(a, b, acc, k_dim=k, barrier=bar, arrive=True)","handlingStrategy":"validation","validationCode":"if arrive:\n    assert getattr(barrier.inner_aval.dtype, 'orders_tensor_core', False)","typeGuard":"def barrier_orders_tensor_core(barrier):\n    return getattr(getattr(barrier, 'inner_aval', barrier).dtype, 'orders_tensor_core', False)","tryCatchPattern":null,"preventionTips":["Create separate barriers for tensor-core-ordered arrivals","Name barriers by purpose (mma vs copy)"],"tags":["jax","pallas","tcgen05","barrier","pipelining"],"backgroundTag":"gpu-barrier-misconfiguration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}