{"record":{"id":"8c4fa3749030da19","repo":"jax-ml/jax","slug":"collective-attribute-is-only-defined-for-tmem-refs","errorCode":null,"errorMessage":"collective attribute is only defined for TMEM refs","messagePattern":"collective attribute is only defined for TMEM refs","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":634,"sourceCode":"\n  def _setitem(self, tracer, idx, value):\n    del tracer, idx, value  # Unused.\n    raise ValueError(\"Ref unions can't be assigned to.\")\n\n  def update(self, inner_aval=None, memory_space=None, kind=None):\n    ref = super().update(inner_aval, memory_space, kind)\n    return AbstractRefUnion(ref.inner_aval, self.refs, self.memory_space)\n\n  @functools.cached_property\n  def layout(self) -> tcgen05.TMEMLayout:\n    if self.memory_space != TMEM:\n      raise ValueError(\"layout attribute is only defined for TMEM refs\")\n    return tcgen05.tmem_default_layout(packing=1)\n\n  @functools.cached_property\n  def collective(self) -> bool:\n    if self.memory_space != TMEM:\n      raise ValueError(\"collective attribute is only defined for TMEM refs\")\n    ref_leaves = jax.tree.leaves(self.refs)\n    first_ref = ref_leaves[0]\n    assert all(ref.collective == first_ref.collective for ref in ref_leaves)\n    return first_ref.collective\n\n  def __eq__(self, other):\n    return (\n        type(self) is type(other)\n        and self.inner_aval == other.inner_aval\n        and self.memory_space == other.memory_space\n        and self.refs == other.refs\n    )\n\n  def __hash__(self):\n    # `flatten_ref_union(self)` creates `TransformedRef`s that refer to `self`,\n    # so we extract the transforms from the `TransformedRef`s in order to avoid\n    # infinite recursion.\n    all_transforms = tuple(","sourceCodeStart":616,"sourceCodeEnd":652,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L616-L652","documentation":"AbstractRefUnion.collective in Mosaic GPU core returns whether the aliased TMEM references use the tcgen05 collective (warpgroup-wide) access mode. It is only defined for TMEM unions; accessing it on SMEM refs raises ValueError. The value is taken from the first leaf ref, asserting all leaves agree.","triggerScenarios":"Accessing `.collective` on an AbstractRefUnion with memory_space != TMEM, e.g. when interpreting or lowering kernels that mix SMEM and TMEM refs (tests like testArgAllReduce / testCommAssocCollective exercise the TMEM path).","commonSituations":"Generic introspection over ref avals during tracing/interpretation; refactors that pass SMEM refs into code paths written for TMEM collective ops; JAX version upgrades where the collective attribute was added to the API surface.","solutions":["Guard with `ref.memory_space == TMEM` before reading `.collective`","When creating TMEM refs, specify collective explicitly (tcgen05 allocation API) so the attribute is well-defined","Audit code paths that assume collective semantics and split TMEM vs SMEM handling"],"exampleFix":"# before\nis_collective = ref_aval.collective\n\n# after\nis_collective = (\n    ref_aval.collective\n    if ref_aval.memory_space == mosaic_gpu_core.TMEM\n    else False\n)","handlingStrategy":"validation","validationCode":"is_collective = (\n    ref_aval.collective\n    if ref_aval.memory_space == mosaic_gpu_core.TMEM\n    else False\n)","typeGuard":"def is_tmem_ref_aval(a) -> bool:\n    return getattr(a, 'memory_space', None) == mosaic_gpu_core.TMEM","tryCatchPattern":"null","preventionTips":["Never read collective on SMEM refs","Set collective explicitly when allocating TMEM refs"],"tags":["jax","pallas","mosaic-gpu","tmem","collective-ops"],"backgroundTag":"invalid-attribute-access","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}