jax-ml/jax · error · ValueError

4-bit block scaled MMA only supports K-fastest operands, but

Error message

4-bit block scaled MMA only supports K-fastest operands, but B is N-fastest

What it means

4-bit block-scaled MMA requires the B operand to be K-fastest as well. If B's fastest-varying dimension is N, mma raises this error.

Source

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

      (b_desc_base, b_k_instr_strides),
      (b_n_group_stride, b_k_group_stride),
      b_fastest,
  ) = mma_utils.create_descriptor(
      b,
      swizzle=b_swizzle,
      group_size=(k_group_elems, n_group_elems),
      logical_k_major=True,
      mma_bytewidth_k=64 if is_sparse else 32,
      split_const=True,
  )

  if is_scaled and utils.bitwidth(mma_a_element_type) == 4:
    if a_fastest != mma_utils.Dim.K:
      raise ValueError(
          "4-bit block scaled MMA only supports K-fastest operands, but A is M-fastest"
      )
    if b_fastest != mma_utils.Dim.K:
      raise ValueError(
          "4-bit block scaled MMA only supports K-fastest operands, but B is N-fastest"
      )
  if is_sparse:
    if b_swizzle == 32 and b_fastest == mma_utils.Dim.K:
      raise NotImplementedError(
          "B tiling too small. Increase swizzle or transpose the input."
      )

  # Step 4. Issue the instructions.
  true = arith.constant(ir.IntegerType.get_signless(1), 1)
  n_collective_group_elems = n_group_elems * num_cta
  n_col_groups = n_groups // n_lane_groups
  assert d.layout.base_tile_shape[0] % 4 == 0
  lanes_per_n_group = d.layout.base_tile_shape[0] // 4
  a_sparse_addr_base = a_sparse_metadata.address if is_sparse else None
  a_scale_addr_base = a_scale.address if is_scaled else None  # pyrefly: ignore[missing-attribute]
  b_scale_addr_base = b_scale.address if is_scaled else None  # pyrefly: ignore[missing-attribute]
  # B scales are padded when N is short, so it can't be derived from n_collective_group_elems.

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Store/transpose B so K is the fastest dimension
2. Adjust b_swizzle to a value compatible with the K-fastest layout

Example fix

# before
b = weights  # N-fastest
tcgen05.mma(a, b, d, a_scale=asc, b_scale=bsc, scale_block=16)
# after
b_t = transpose_to_k_fastest(weights)
tcgen05.mma(a, b_t, d, a_scale=asc, b_scale=bsc, scale_block=16)

assert b_fastest == mma_utils.Dim.K, '4-bit scaled MMA needs K-fastest B'

When it happens

Trigger: Calling mma() with is_scaled=True, 4-bit operands, and b_fastest == Dim.N (B stored row-major over (N, K) transposed).

Common situations: Standard GEMM weight layout (N-major B) reused from an 8-bit kernel; missing transpose of the weights for NVFP4.

Related errors

• 4-bit block scaled MMA only supports K-fastest operands, but
• Expected B scales to have a M=128 layout, got {b_scale.layou
• MMA with element type {elem_type_str} does not support block
• Scale element type mismatch: expected f8e8m0fnu or f8e4m3fn,
• Unsupported element type for block scaling: {a_element_type}

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