jax-ml/jax · error · ValueError

Unsupported A register array layout

Error message

Unsupported A register array layout

What it means

At wgmma.py:164, an in-register A FragmentedArray must be laid out as fa.WGMMA_LAYOUT (or fa.WGMMA_LAYOUT_8BIT for byte types) so registers match the wgmma A-operand fragment layout.

Source

Thrown at jax/experimental/mosaic/gpu/wgmma.py:164

  f8e5m2 = ir.Float8E5M2Type.get()
  f8e4m3fn = ir.Float8E4M3FNType.get()
  if b_k_stride % 16:
    raise ValueError
  assert bytewidth(a_element_type) == bytewidth(b_element_type)
  # Only 16-bit types support transposes
  supports_transpose = bytewidth(b_element_type) == 2
  if not supports_transpose and (a_transpose or b_transpose):
    raise ValueError("Only f16 WGMMA supports transposes")
  if a_in_regs := isinstance(a, fa.FragmentedArray):
    if a.mlir_dtype not in {bf16, f16, i8, f8e5m2, f8e4m3fn}:
      raise ValueError(f"Unsupported A register array dtype: {a.mlir_dtype}")
    # Column count must be equal to swizzle // bytewidth.
    elt_bytewidth = utils.bytewidth(a_element_type)
    swizzle_elems = swizzle // elt_bytewidth
    if a.shape != (64, swizzle_elems):
      raise ValueError("Unsupported A register array shape")
    if a.layout not in {fa.WGMMA_LAYOUT, fa.WGMMA_LAYOUT_8BIT}:
      raise ValueError("Unsupported A register array layout")
    if a_k_stride is not None or a_transpose is not None:
      raise ValueError("Unsupported WGMMA features with A in registers")
  else:
    if a_k_stride is None or a_k_stride % 16:
      raise ValueError
    if a_transpose is None:
      raise ValueError

  if isinstance(out_ty, ir.F32Type) or out_ty == i32:
    num_acc_regs = n // 2
    out_ty_field = ir.VectorType.get((1,), out_ty)
    acc_regs = list(acc.flat)
    assert acc_regs[0].type == ir.VectorType.get((1,), out_ty)
    to_acc_vec_regs = lambda regs: np.array(regs).reshape(acc.shape)
    acc_constraint = "r" if isinstance(out_ty, ir.IntegerType) else "f"
  elif isinstance(out_ty, ir.F16Type):
    num_acc_regs = n // 4
    out_ty_field = i32

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Convert: a = a.to_layout(fa.WGMMA_LAYOUT_8BIT if 8-bit else fa.WGMMA_LAYOUT)
2. Pass A from SMEM instead of registers when the layout cannot be preserved
3. Verify layout constants exist in your installed JAX/mosaic version

Example fix

# before
acc = wgmma.wgmma(a_rowmajor, b, acc, ...)
# after
a = a_rowmajor.to_layout(fa.WGMMA_LAYOUT)
acc = wgmma.wgmma(a, b, acc, ...)

assert a.layout in {fa.WGMMA_LAYOUT, fa.WGMMA_LAYOUT_8BIT}, 'convert A layout first'

def wgmma_a_layout_ok(a): return a.layout in {fa.WGMMA_LAYOUT, fa.WGMMA_LAYOUT_8BIT}

When it happens

Trigger: Passing a as a FragmentedArray in a row-major or blocked layout when using the A-in-registers form of wgmma.

Common situations: Reusing a register tensor produced by earlier elementwise ops or a previous wgmma result without relayout; layout constant renames across JAX versions.

Related errors

• Only WGMMA layouts supported in WGMMAAccumulator
• Source and destination layouts aren't compatible for a broad
• packed cannot be specified if layout is specified.
• Cannot commute `UntilingTransform` with `ReshapeTransform` w
• Swizzle {self.swizzle} requires the trailing dimension to be

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