jax-ml/jax · error · ValueError

Either none or both scales should be provided

Error message

Either none or both scales should be provided

What it means

Block-scaled MMA requires scale factors for both operands or neither — the hardware descriptor has symmetric A/B scale fields. Passing only a_scale (or only b_scale) is an API misuse and raises ValueError.

Source

Thrown at jax/experimental/mosaic/gpu/tcgen05.py:200

    b: ir.Value,
    *,
    a_swizzle: int = 128,
    b_swizzle: int = 128,
    a_scale: TMEMRef | None = None,
    b_scale: TMEMRef | None = None,
    a_sparse_metadata: TMEMRef | None = None,
    accumulate: ir.Value | bool = True,
    collective: bool = False,
) -> None:
  if a_swizzle == 16 or b_swizzle == 16:
    raise NotImplementedError("No swizzle is not supported")
  i8 = ir.IntegerType.get_signless(8)
  i32 = ir.IntegerType.get_signless(32)
  if isinstance(accumulate, bool):
    accumulate = arith.constant(ir.IntegerType.get_signless(1), accumulate)
  num_cta = 2 if collective else 1
  if (is_scaled := a_scale is not None) != (b_scale is not None):
    raise ValueError("Either none or both scales should be provided")
  is_sparse = a_sparse_metadata is not None
  if is_scaled and is_sparse:
    if isinstance(a, TMEMRef):
      raise NotImplementedError(
          "A in TMEM unsupported for block-scaled sparse matmuls"
      )

  # Step 1. Establish the shape and element type of the operation.
  if not isinstance(b.type, ir.MemRefType):
    raise ValueError(f"B must be a memref, got: {b.type}")
  (k, n), b_element_type = mma_utils.tiled_memref_shape(b)
  if isinstance(a, TMEMRef):
    m, k2 = a.shape
    a_element_type = a.dtype
    if m != 128:
      raise NotImplementedError(f"Only M=128 is supported for MMA with A in TMEM, but got M={m}")
    # Watch out: this layout must be consistent with D's layout (up to packing).
    expected_packing = 32 // utils.bitwidth(a_element_type)

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Pass both a_scale and b_scale, or omit both
2. If only one operand needs scaling, pass an identity/1.0-filled scale for the other operand

Example fix

# before
mma(acc, a, b, a_scale=a_scale)

# after
mma(acc, a, b, a_scale=a_scale, b_scale=b_scale)  # or drop both

if (a_scale is None) != (b_scale is None):
    raise ValueError('provide both a_scale and b_scale, or neither')
mma(acc, a, b, a_scale=a_scale, b_scale=b_scale)

When it happens

Trigger: Calling mma(..., a_scale=tref) without b_scale, or vice versa; conditionally passing one scale based on a None check.

Common situations: Refactoring code that had scales on only one operand in a different API; optional-argument plumbing where one scale fails to propagate.

Related errors

• Unsupported scale type: {scale_type}
• multi_dot requires at least two arrays; got len(arrays)={len
• a_scale and b_scale must both be present or absent.
• Expected both or neither of scales to be specified.
• Unsupported accumulator dtype: {acc_dtype}

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/426322a4d8aee17c. Report an issue: GitHub.