jax-ml/jax · error · NotImplementedError

Unhandled transforms for multimem_load_reduce: {transforms}

Error message

Unhandled transforms for multimem_load_reduce: {transforms}

What it means

In the Lane-semantics lowering of multimem_load_reduce, _handle_transforms was invoked with allow_peer_refs=False; if any transforms remain on the ref after handling (e.g. peer-ref transforms), they are unsupported and the lowering aborts.

Source

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

    raise RuntimeError(
        "Failed to infer the output layout of multimem_load_reduce. Please apply"
        " plgpu.layout_cast to its output right after its creation."
    )
  if not isinstance(layout, (mgpu.TiledLayout, mgpu.WGStridedFragLayout)):
    raise ValueError(
        "Only tiled and WG strided layouts are supported by"
        f" multimem_load_reduce, but got {layout}"
    )
  dtype = ctx.avals_out[0].dtype
  transforms = tree.unflatten(transforms_leaves)
  transform_avals = tree.unflatten(ctx.avals_in[1:])
  ref_aval = ctx.avals_in[0]
  assert isinstance(ref_aval, state_types.AbstractRef)
  ref, _, transforms = lowering._handle_transforms(ctx, ref_aval, ref,
                                                   transform_avals, transforms,
                                                   allow_peer_refs=False)
  if transforms:
    raise NotImplementedError(
        f"Unhandled transforms for multimem_load_reduce: {transforms}"
    )
  multi_ref = ctx.launch_ctx.to_remote_multicast(ref)
  is_signed = mgpu_utils.is_signed(dtype)
  arr = mgpu.FragmentedArray.load_reduce_untiled(
      multi_ref, layout=layout, is_signed=is_signed, reduction=reduction_op
  )
  return arr


@lowering.register_lowering_rule(multimem_load_reduce_p, mgpu.LoweringSemantics.Warpgroup)
def _multimem_load_reduce_lowering_rule_wg(
    ctx: lowering.LoweringRuleContext, ref, *transforms_leaves, tree, collective_axes, reduction_op,
):
  if (mesh_info := ctx.module_ctx.mesh_info) is None:
    raise ValueError(
        "JAX device mesh is required by multimem_load_reduce, but not defined."
    )

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

Use a locally-created ref (allocated in this kernel) rather than a peer ref
Strip/avoid extra transforms on the ref before the call
Update JAX to pick up broader transform coverage in _handle_transforms
Fall back to a normal load plus an explicit collective reduce

Defensive patterns

Strategy: fallback

Try / catch

try:
    kernel_jit(x)
except NotImplementedError as e:
    if 'multimem_load_reduce' in str(e):
        run_load_plus_psum_fallback(x)

Prevention

Avoid peer refs and stacked transforms on multimem inputs
Maintain a load+reduce fallback path for transform edge cases

When it happens

Trigger: Passing a ref with residual or peer-related transforms to multimem_load_reduce in a kernel lowered under Lane semantics.

Common situations: Cross-shard ref sharing feeding a multimem load; combining multimem ops with experimental ref-transform utilities; regressions after JAX upgrades changing transform semantics.

Related errors

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