{"record":{"id":"ab16150f7e1be4df","repo":"jax-ml/jax","slug":"some-aliased-tmem-references-are-collective-and-so","errorCode":null,"errorMessage":"Some aliased TMEM references are collective and some are not.","messagePattern":"Some aliased TMEM references are collective and some are not\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":690,"sourceCode":"\n  def __init__(self, *refs: _GPUMemoryRefTree):\n    ref_leaves = jax.tree.leaves(refs)\n    if all(ref.memory_space == SMEM for ref in ref_leaves):\n      object.__setattr__(self, \"refs\", refs)\n      num_bytes = max(map(_ref_group_size, self.refs))\n      super().__init__(\n          inner_aval=jax_core.ShapedArray(\n              (num_bytes,), jnp.int8\n          ),\n          memory_space=SMEM,\n          transforms=(),\n      )\n    elif all(ref.memory_space == TMEM for ref in ref_leaves):\n      object.__setattr__(self, \"refs\", refs)\n      max_cols = max(map(_ref_group_tmem_col_size, self.refs))\n      is_collective = ref_leaves[0].collective\n      if any(r.collective != is_collective for r in ref_leaves):\n        raise ValueError(\n            \"Some aliased TMEM references are collective and some are not.\"\n        )\n      super().__init__(\n          inner_aval=jax_core.ShapedArray(\n              shape=(128, max_cols,),\n              dtype=jnp.int32,\n          ),\n          memory_space=TMEM,\n          transforms=(),\n          layout=tcgen05.tmem_default_layout(packing=1),\n          collective=all(ref.collective for ref in ref_leaves),\n      )\n    else:\n      raise NotImplementedError(\n          \"All aliased Refs must have the same memory space (SMEM or TMEM). \"\n          f\"Got {(ref.memory_space for ref in ref_leaves)}.\")\n\n  def get_ref_aval(self) -> AbstractRefUnion:","sourceCodeStart":672,"sourceCodeEnd":708,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L672-L708","documentation":"RefMap (a union of aliased refs) in Mosaic GPU requires that when all aliased refs live in TMEM, they must all have the same `collective` flag, since the union exposes a single collective value. Mixing collective and non-collective TMEM refs is ambiguous and raises ValueError at construction.","triggerScenarios":"Constructing a RefMap (RefUnion) whose pytree leaves include TMEM refs created with different collective settings, e.g. combining refs allocated via `tcgen05.alloc` collective API with per-thread TMEM refs.","commonSituations":"Refactoring kernels to share/alias TMEM buffers between collective matmul ops and per-warpgroup ops; version upgrades that introduced the collective invariant; tests constructing mixed ref trees for aliasing analysis.","solutions":["Create all TMEM refs that will be aliased with the same collective mode (pass matching collective= argument to the tcgen05 allocation)","Split the RefMap so collective and non-collective refs are not aliased together","Check `ref.collective` on each TMEM ref before building the RefMap and adjust allocation"],"exampleFix":"# before\nref_a = tcgen05.alloc(..., collective=True)\nref_b = tcgen05.alloc(..., collective=False)\nref_map = RefMap((ref_a, ref_b))  # ValueError\n\n# after\nref_b = tcgen05.alloc(..., collective=True)\nref_map = RefMap((ref_a, ref_b))","handlingStrategy":"validation","validationCode":"leaves = jax.tree.leaves(refs)\nassert all(\n    r.memory_space != mosaic_gpu_core.TMEM\n    or r.collective == leaves[0].collective\n    for r in leaves\n), 'mixed collective modes in aliased TMEM refs'","typeGuard":"def refs_have_uniform_collective(refs) -> bool:\n    leaves = jax.tree.leaves(refs)\n    tmem = [r for r in leaves if r.memory_space == mosaic_gpu_core.TMEM]\n    return len({r.collective for r in tmem}) <= 1","tryCatchPattern":"null","preventionTips":["Allocate all aliased TMEM refs with the same collective flag","Group refs by (memory_space, collective) before building RefMaps"],"tags":["jax","pallas","mosaic-gpu","tmem","aliasing"],"backgroundTag":"invariant-violation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}