{"record":{"id":"a8f4d6307132df7f","repo":"jax-ml/jax","slug":"ref-must-be-a-reference","errorCode":null,"errorMessage":"ref must be a reference","messagePattern":"ref must be a reference","errorType":"exception","errorClass":"TypeError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":1053,"sourceCode":"\n  def undo(self, x: jax_core.AbstractValue) -> state_types.Transform:\n    raise NotImplementedError()\n\n  def commute_ndindexer(\n      self, _: jax_core.AbstractValue, indexer: indexing.NDIndexer\n  ) -> tuple[indexing.NDIndexer, MulticastRef]:\n    return indexer, self\n\n\ndef remote_ref(\n    ref: _Ref,\n    device_id: jax.typing.ArrayLike,\n    device_id_type: pallas_primitives.DeviceIdType = pallas_primitives.DeviceIdType.MESH,\n) -> pallas_core.TransformedRef:\n  \"\"\"Translate memref to a symmetric memref on a peer device.\"\"\"\n  if not isinstance(ref, pallas_core.TransformedRef):\n    if not isinstance(jax_core.typeof(ref), state_types.AbstractRef):\n      raise TypeError(\"ref must be a reference\")\n    ref = pallas_core.TransformedRef(ref, transforms=())\n  if any(isinstance(t, MulticastRef) for t in ref.transforms):\n    raise ValueError(\"Can't make a multicast reference into a peer reference.\")\n  return pallas_core.TransformedRef(\n      ref.ref, (*ref.transforms, PeerMemRef(device_id, device_id_type)),\n  )\n\n\n@tree_util.register_dataclass\n@dataclasses.dataclass(frozen=True)\nclass ClusterRefTransform(state_types.Transform):\n  dims: tuple[jax_core.AxisName, ...] = jax.tree.static()\n  idxs: tuple[Any, ...]\n\n  def __post_init__(self):\n    if len(self.dims) != len(self.idxs):\n      raise ValueError(\"dims and idxs must have the same length\")\n","sourceCodeStart":1035,"sourceCodeEnd":1071,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L1035-L1071","documentation":"remote_ref requires its argument to be a reference (an AbstractRef or TransformedRef). Passing a plain array, tracer, or other value raises TypeError('ref must be a reference').","triggerScenarios":"Calling pallas_mosaic_gpu.remote_ref(array, device_id) with a non-ref value, e.g. passing a block from a kernel argument that isn't a Ref, or passing an already-consumed value.","commonSituations":"Writing cross-device TPU kernels and passing jax arrays or scalars where a memref is expected; passing a ref created by a different pallas backend.","solutions":["Pass a pallas Ref / TransformedRef obtained from the kernel's input/output parameters","Ensure the value comes from the kernel signature (Ref[...]) rather than a computed array","Wrap raw refs consistently: let remote_ref wrap plain AbstractRefs itself"],"exampleFix":"// before\npeer = remote_ref(some_array, peer_id)\n\n// after\npeer = remote_ref(out_ref, peer_id)  # out_ref is a Ref from kernel params","handlingStrategy":"type-guard","validationCode":"from jax._src.pallas import state_types\nassert isinstance(ref, pallas_core.TransformedRef) or isinstance(jax_core.typeof(ref), state_types.AbstractRef)","typeGuard":"def is_ref(x): return isinstance(x, pallas_core.TransformedRef) or type(jax_core.typeof(x)).__name__ == 'AbstractRef'","tryCatchPattern":null,"preventionTips":["Only pass kernel Ref parameters to remote_ref"],"tags":["jax","pallas","mosaic-gpu","peer-memory","type-error"],"backgroundTag":"invalid-argument-type","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}