{"record":{"id":"cdeb138751335de0","repo":"jax-ml/jax","slug":"expected-both-or-neither-of-scales-to-be-specified","errorCode":null,"errorMessage":"Expected both or neither of scales to be specified.","messagePattern":"Expected both or neither of scales to be specified\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/dialect_lowering.py","lineNumber":2515,"sourceCode":") -> Sequence[ir.Value]:\n  ctx.check_collective(op)\n\n  def tmem_layout(operand):\n    result = inference_utils.in_tmem_layout_for_operand(op, operand)\n    # satisfy the type checker\n    assert result is not None\n    return result\n\n  acc_ref = _tmem_ref_from_ir(op.accumulator, tmem_layout(op.accumulator))\n\n  if op.a_sparse_metadata is not None:\n    a_sparse_metadata = _tmem_ref_from_ir(\n        op.a_sparse_metadata, tmem_layout(op.a_sparse_metadata)\n    )\n  else:\n    a_sparse_metadata = None\n  if (scaled := op.a_scale is not None) != (op.b_scale is not None):\n    raise ValueError(\"Expected both or neither of scales to be specified.\")\n  if scaled:\n    a_scale = _tmem_ref_from_ir(op.a_scale, tmem_layout(op.a_scale))  # pyrefly: ignore[bad-argument-type]\n    b_scale = _tmem_ref_from_ir(op.b_scale, tmem_layout(op.b_scale))  # pyrefly: ignore[bad-argument-type]\n  else:\n    a_scale = None\n    b_scale = None\n\n  if utils.is_smem_ref(op.a):\n    a_transforms, b_transforms = inference_utils.in_transforms(op)\n    a_swizzle = swizzle_from_transforms_attr(a_transforms)\n    b_swizzle = swizzle_from_transforms_attr(b_transforms)\n    a_ref = unwrap_transformed_memref(op.a, a_transforms)\n    b_ref = unwrap_transformed_memref(op.b, b_transforms)\n  else:\n    a_ref = _tmem_ref_from_ir(op.a, tmem_layout(op.a))\n    [b_transforms] = inference_utils.in_transforms(op)\n    b_swizzle = swizzle_from_transforms_attr(b_transforms)\n    a_swizzle = b_swizzle","sourceCodeStart":2497,"sourceCodeEnd":2533,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/dialect_lowering.py#L2497-L2533","documentation":"tcgen05 MMA lowering requires that a_scale and b_scale are either both provided (scaled MMA, e.g. FP8 with MX formats) or both omitted. Supplying only one is ambiguous and rejected.","triggerScenarios":"Calling tcgen05_mma with a_scale set but b_scale None, or vice versa — typically when building scaled matmuls with per-operand scale tensors.","commonSituations":"FP8/MXFP8 matmul kernels where the B operand's scale is accidentally dropped (e.g. None default from a wrapper) or where a user assumes scales are per-operand optional.","solutions":["Pass both a_scale and b_scale, or neither","Audit wrapper functions for default None scale arguments leaking through","If only one side needs scaling, use scale=1 explicit constant scale tensor for the other side if the op semantics allow"],"exampleFix":"// before\nmgpu.tcgen05_mma(a, b, acc, a_scale=sa)  # b_scale missing\n// after\nmgpu.tcgen05_mma(a, b, acc, a_scale=sa, b_scale=sb)","handlingStrategy":"validation","validationCode":"assert (a_scale is None) == (b_scale is None), 'provide both scales or neither'","typeGuard":"def valid_scales(a_scale, b_scale) -> bool:\n    return (a_scale is None) == (b_scale is None)","tryCatchPattern":null,"preventionTips":["Wrap tcgen05_mma in a helper enforcing the paired-scales invariant","Treat scales as a single optional tuple argument in your API"],"tags":["mosaic","gpu","tcgen05","mma","scales","argument-validation"],"backgroundTag":"mismatched-paired-arguments","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}