{"record":{"id":"7a300b024b30387b","repo":"jax-ml/jax","slug":"swizzle-self-swizzle-is-not-supported-only-32","errorCode":null,"errorMessage":"Swizzle {self.swizzle} is not supported. Only 32, 64 and 128 are accepted.","messagePattern":"Swizzle (.+?) is not supported\\. Only 32, 64 and 128 are accepted\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":1173,"sourceCode":"    return cls(dtypes.dtype(ref.dtype), ref.ref.shape, byte_offset, alias_group_idx, layout)\n\n  def transform_type(self, x):\n    match x:\n      case state_types.AbstractRef():\n        return x.update(inner_aval=self.transform_type(x.inner_aval))\n      case jax_core.ShapedArray():\n        return x.update(shape=self.shape, dtype=self.dtype)\n      case _:\n        raise TypeError(f\"Unsupported type: {x}\")\n\n\n@dataclasses.dataclass(frozen=True)\nclass SwizzleTransform(state_types.Transform):\n  swizzle: int\n\n  def __post_init__(self):\n    if self.swizzle not in {32, 64, 128}:\n      raise ValueError(\n          f\"Swizzle {self.swizzle} is not supported. Only 32, 64 and 128 are\"\n          \" accepted.\"\n      )\n\n  def transform_type(\n      self, x: jax_core.AbstractValue\n  ) -> jax_core.AbstractValue:\n    match x:\n      case jax_core.ShapedArray():\n        swizzle_elems = (self.swizzle * 8) // dtypes.itemsize_bits(x.dtype)\n        if swizzle_elems != x.shape[-1]:\n          raise ValueError(\n              f\"Swizzle {self.swizzle} requires the trailing dimension to be of\"\n              f\" size {swizzle_elems}, but got shape: {x.shape}\"\n          )\n        return x\n      case state_types.AbstractRef():\n        return x.update(inner_aval=self.transform_type(x.inner_aval))","sourceCodeStart":1155,"sourceCodeEnd":1191,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L1155-L1191","documentation":"SwizzleTransform only accepts swizzle values 32, 64, or 128 (byte widths understood by TPU shared-memory swizzling hardware); anything else fails validation in __post_init__.","triggerScenarios":"Constructing SwizzleTransform(swizzle=16) or other invalid values, usually by parameterizing swizzle from user config in a pallas kernel.","commonSituations":"Passing a swizzle derived from dtype size or block shape arithmetic that lands outside {32, 64, 128}; copying example code with a wrong constant.","solutions":["Clamp/choose swizzle from {32, 64, 128} only","Derive swizzle from the TMA block layout (typically 128) rather than computing arbitrary values","Validate user-supplied swizzle before constructing the transform"],"exampleFix":"# before\nSwizzleTransform(swizzle=16)\n\n# after\nassert swizzle in (32, 64, 128)\nSwizzleTransform(swizzle=swizzle)","handlingStrategy":"validation","validationCode":"if swizzle not in (32, 64, 128): raise ValueError('swizzle must be 32, 64 or 128')","typeGuard":"def valid_swizzle(s): return s in (32, 64, 128)","tryCatchPattern":null,"preventionTips":["Hard-code swizzle choices; validate config inputs early"],"tags":["jax","pallas","mosaic-gpu","swizzle","validation"],"backgroundTag":"invalid-config-value","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}