jax-ml/jax · error · ValueError

Unsupported load reduce operation: {orig_reduce}

Error message

Unsupported load reduce operation: {orig_reduce}

What it means

The reduce argument of async_load_tmem must be one of the supported reduction names handled by the match statement (min/max variants including absmin and absmax). An unknown value raises ValueError with the original op.

Source

Thrown at jax/_src/pallas/mosaic_gpu/primitives.py:4166

        f"Unimplemented transforms for TMEM refs. {transforms=}"
    )
  if reduce is not None:
    orig_reduce = reduce
    # Make sure we will interpret the reduction as unsigned, since we erase
    # signedness.
    if isinstance(x_aval.dtype, jnp.unsignedinteger) and "abs" not in reduce:
      reduce = "abs" + reduce  # type: ignore
    match reduce:
      case "min":
        reduce_attr = mgpu.dialect.TMEMLoadReduction.Min  # pyrefly: ignore[missing-attribute]
      case "max":
        reduce_attr = mgpu.dialect.TMEMLoadReduction.Max  # pyrefly: ignore[missing-attribute]
      case "absmin":
        reduce_attr = mgpu.dialect.TMEMLoadReduction.AbsMin  # pyrefly: ignore[missing-attribute]
      case "absmax":
        reduce_attr = mgpu.dialect.TMEMLoadReduction.AbsMax  # pyrefly: ignore[missing-attribute]
      case _:
        raise ValueError(f"Unsupported load reduce operation: {orig_reduce}")
    return tuple(mgpu.dialect.async_load_tmem(x_tmem, reduce=reduce_attr))  # type: ignore
  return (mgpu.dialect.async_load_tmem(x_tmem),)


wait_load_tmem_p = jax_core.Primitive("wait_load_tmem")
wait_load_tmem_p.multiple_results = True

def wait_load_tmem():
  """Awaits all previously asynchronous TMEM loads issued by the calling thread.

  Once this function returns, the TMEM loads issued by the calling thread are
  guaranteed to have completed. The read TMEM regions can be safely overwritten
  by the calling thread, or any threads signalled through ``Barrier``s with
  ``orders_tensor_core=True``.
  """
  wait_load_tmem_p.bind()

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Use one of the supported names: min, max, absmin, absmax (per tcgen05 TMEMLoadReduction)
2. Upgrade JAX if you expect a newly added reduction op

Example fix

// before
v, r = async_load_tmem(ref, reduce='minimum')
// after
v, r = async_load_tmem(ref, reduce='min')

assert reduce in ('min','max','absmin','absmax'), reduce

def valid_reduce(op): return op in {'min','max','absmin','absmax'}

When it happens

Trigger: Passing reduce='mean' or a misspelled or too-new reduction name to async_load_tmem.

Common situations: Typos; using reduction names from a different API (e.g. jnp reduction names) or a newer JAX version's names on an older install.

Understand the failure class

Background: Invalid enum value errors: "Unknown type", "Invalid scope", "must be one of" — when a string is not on the library's allowed list — this error's family across 23 libraries.

Related errors

• Unsupported dtype for reduction: {val_aval.dtype}. Only floa
• Reductions over unsigned integers not implemented.
• Reductions over {x_aval.dtype} not implemented.
• Only single axis reduction supported
• reductions require axes to be (0,) on SparseCore, but got {a

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