jax-ml/jax · error · NotImplementedError

Only SMEM and TMEM refs are supported.

Error message

Only SMEM and TMEM refs are supported.

What it means

`flatten_ref_union` only knows how to lay out refs living in shared memory (SMEM) or tensor memory (TMEM). Calling it on refs whose memory space is anything else — typically global memory (GMEM) refs or host/abstract refs — raises NotImplementedError.

Source

Thrown at jax/_src/pallas/mosaic_gpu/core.py:595

      def unflatten(ref):
        nonlocal col_offset
        col_offset = align_to(col_offset, TMEM_COL_ALIGNMENT)
        if not isinstance(ref, pallas_core.TransformedRef):
          ref = pallas_core.TransformedRef(ref, transforms=())
        ncols = ref.layout.cols_in_shape(ref.shape,
                                         dtypes.itemsize_bits(ref.dtype))
        transform = ExtractAliasedRef.from_transformed_ref(
            ref, col_offset, group_idx, layout=ref.layout)
        result = pallas_core.TransformedRef(
            ref_union, transforms=(transform, *ref.transforms)
        )
        col_offset += ncols
        return result
      flat_refs.append(jax.tree.map(unflatten, ref_group))
      union_cols = max(union_cols, col_offset)
    assert union_cols == ref_union.shape[1], (union_cols, ref_union.shape[1])
  else:
    raise NotImplementedError("Only SMEM and TMEM refs are supported.")
  return tuple(flat_refs)


class AbstractRefUnion(state.AbstractRef):
  refs: Sequence[_GPUMemoryRefTree]

  def __init__(
      self,
      aval,
      refs: Sequence[_GPUMemoryRefTree],
      memory_space,
  ):
    self.refs = refs
    super().__init__(aval, memory_space=memory_space)

  def _iter(self, tracer):
    return iter(flatten_ref_union(tracer))

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Filter the ref tree so only SMEM/TMEM scratch refs reach flatten_ref_union
2. Keep GMEM input/output refs out of the union; process them separately in the kernel
3. Check ref.memory_space before grouping if refs come from mixed sources

Example fix

# before
flat = flatten_ref_union((x_ref, smem_scratch))  # x_ref is a GMEM input
# after
flat = flatten_ref_union((smem_scratch,))  # only SMEM/TMEM refs

def only_onchip(refs):
    return [r for r in refs if type(r.memory_space).__name__ in ('SMEMAddressSpace','TMEMAddressSpace') or str(r.memory_space) in ('SMEM','TMEM')]

def is_onchip_ref(ref) -> bool:
    ms = getattr(ref, 'memory_space', None)
    return ms is not None and str(getattr(ms, 'memory_space', ms)) in ('SMEM', 'TMEM')

When it happens

Trigger: Passing a grid/gmem BlockRef (an input or output of a Pallas kernel) into `flatten_ref_union`, or any ref tree containing a ref whose memory_space is not SMEM or TMEM.

Common situations: Accidentally feeding kernel I/O refs (which live in GMEM) into union-building helpers meant for scratch buffers; generic tree-mapping code that hits both scratch and I/O refs; assuming the helper is memory-space agnostic.

Related errors

• Unsupported memory space.
• Unsupported memory space: {x.memory_space}
• Unsupported dtype: {ref.dtype}
• Ref unions can't be assigned to.
• Unsupported transform: {type(transform)}

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