jax-ml/jax · error · ValueError

Unsupported reduction kind: {reduce}

Error message

Unsupported reduction kind: {reduce}

What it means

In _load_32xcols_native's fused reduction loop, floating-point reductions only implement min/absmin (minimumf) and max/absmax (maximumf). Any other reduce string that survived earlier validation reaches the match's default branch and raises ValueError.

Source

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

  regs = [None] * (cols // reg_packing)
  red_reg = None
  for addr_row_col, instr_num, lane_step, num_slice in it:
    assert lane_step == 0, lane_step
    instr_regs = _tmem_load(addr_row_col, load_shape, instr_num, pack, reduce, dtype)
    if reduce:
      *instr_regs, instr_red_reg = instr_regs
      instr_red_reg = utils.bitcast(instr_red_reg, dtype)
      if red_reg is None:
        red_reg = instr_red_reg
      elif isinstance(dtype, ir.F32Type):
        # abs is applied by the TMEM load, so both red_regs are non-negative.
        match reduce:
          case "min" | "absmin":
            red_reg = arith.minimumf(red_reg, instr_red_reg)
          case "max" | "absmax":
            red_reg = arith.maximumf(red_reg, instr_red_reg)
          case _:
            raise ValueError(f"Unsupported reduction kind: {reduce}")
      elif isinstance(dtype, ir.IntegerType):
        match reduce:
          case "min":
            red_reg = arith.minsi(red_reg, instr_red_reg)
          case "max":
            red_reg = arith.maxsi(red_reg, instr_red_reg)
          case "absmin":
            red_reg = arith.minui(red_reg, instr_red_reg)
          case "absmax":
            red_reg = arith.maxui(red_reg, instr_red_reg)
          case _:
            raise ValueError(f"Unsupported reduction kind: {reduce}")
      else:
        raise ValueError(f"Unsupported reduction dtype: {dtype}")
    if reg_packing == 1 and vector_length == 2:
      regs[num_slice] = [llvm.bitcast(dtype, r) for r in instr_regs]
    else:
      regs[num_slice] = [utils.bitcast(r, vec_ty) for r in instr_regs]

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Use only min, max, absmin, absmax for float load-reduces
2. If extending Mosaic, add the case to the match (e.g. arith.addf for 'sum') and update the outer validation
3. Call the public TensorMem.load rather than the private helper so validation runs

def is_native_reduce(reduce: str) -> bool:
    return reduce in ('min','max','absmin','absmax')

try:
    arr, red = tmem.load(layout, reduce=reduce)
except ValueError as e:
    if 'Unsupported reduction kind' in str(e):
        arr, _ = tmem.load(layout); red = manual_reduce(arr)
    else: raise

When it happens

Trigger: Internal: calling _load_32xcols_native with a reduce kind other than the float-supported set for an F32 dtype — normally guarded by TensorMem.load's earlier checks, so hitting this means a new/unvalidated reduce kind was passed through.

Common situations: Adding a new reduction kind (e.g. 'sum') to Mosaic and wiring it past the outer validation; calling the private helper directly in custom lowering code.

Related errors

• Unsupported reduction for f32. Only min, max, absmin, and ab
• Unsupported dtype for reduction: {self.dtype}
• Fused load-reduce is not supported for this layout
• Unsupported reduction kind: {op.kind}
• Output layout {out_layout} must match the accumulator layout

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