{"record":{"id":"4e9400faeacfd971","repo":"jax-ml/jax","slug":"no-valid-out-swizzle-what-minor-dimension-has-m","errorCode":null,"errorMessage":"No valid out swizzle{what}: minor dimension has {minor_dim_bits} bits, which is not a multiple of 128 (16 bytes)","messagePattern":"No valid out swizzle(.+?): minor dimension has (.+?) bits, which is not a multiple of 128 \\(16 bytes\\)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/helpers.py","lineNumber":249,"sourceCode":"  applied as ``plgpu`` transform:\n\n    transforms = (\n        plgpu.TilingTransform((8, 8 * swizzle // elem_bits)),\n        plgpu.SwizzleTransform(swizzle))\n    )\n\n  Args:\n    minor_dim_bits: The number of bits in the minor (last) dimension of the\n      memory region. Usually computed as ``dim_size * jnp.finfo(dtype).bits``.\n    what: A string describing the operand for which the swizzle is being\n      computed. Improves the error message if specified.\n  \"\"\"\n  for swizzle_bytes in (128, 64, 32, 16):\n    if minor_dim_bits % (swizzle_bytes * 8) == 0:\n      return swizzle_bytes\n  if what:\n    what = \" for \" + what\n  raise ValueError(\n      f\"No valid out swizzle{what}: minor dimension has\"\n      f\" {minor_dim_bits} bits, which is not a multiple of 128 (16 bytes)\"\n  )\n\n\ndef planar_snake(\n    lin_idx: jax.Array,\n    shape: tuple[int | jax.Array, int | jax.Array],\n    minor_dim: int,\n    tile_width: int,\n):\n  \"\"\"Converts a linear index into an index into shape, trying to optimize locality.\n\n  The \"space filling curve\" this function computes splits the minor dimension\n  into tiles of length ``tile_width``. Every other tile has its major dimension\n  inverted, so that the iteration order \"snakes around\" when going from one tile\n  to another.\n","sourceCodeStart":231,"sourceCodeEnd":267,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/helpers.py#L231-L267","documentation":"find_swizzle picks the largest valid shared-memory swizzle (128/64/32/16 bytes) dividing the minor dimension's bit width. If minor_dim_bits is not a multiple of 128 bits (16 bytes), no swizzle is valid and the error names the offending dimension.","triggerScenarios":"Calling find_swizzle(minor_dim_bits) (directly or via the fused matmul helpers matmul0..matmul5) with a minor dimension whose bit width is not a multiple of 128 — e.g. a 96-bit row, or dtype*elements combos like 8 floats of 12-bit packed types.","commonSituations":"Choosing block shapes/dtypes in a Pallas GPU matmul so that K or N * itemsize_bits isn't a multiple of 128; using 3- or 6-bit packed types or odd vector lengths with the tcgen05/WGMMA path.","solutions":["Adjust the tile's minor dimension (dtype bits x element count) to a multiple of 128 bits, e.g. ensure N * itemsize_bits % 128 == 0","Use a dtype with a standard bit width (8/16/32 bits) and a matching vector length (e.g. 16 elems of 8-bit = 128 bits)","Pass the `what` description string for a clearer error identifying which operand failed"],"exampleFix":"# before: 6 elements of 16-bit = 96 bits\nfind_swizzle(6 * 16)\n\n# after: 8 elements of 16-bit = 128 bits\nfind_swizzle(8 * 16)","handlingStrategy":"validation","validationCode":"assert minor_dim_bits % 128 == 0, f'{minor_dim_bits=} not a multiple of 128 bits (16 bytes)'","typeGuard":"def swizzle_possible(minor_dim_bits: int) -> bool:\n    return minor_dim_bits % 128 == 0 and minor_dim_bits > 0","tryCatchPattern":"try:\n    sw = find_swizzle(minor_dim_bits, what='operand A')\nexcept ValueError:\n    # round the vector length up to a 128-bit multiple\n    minor_dim_bits = -(-minor_dim_bits // 128) * 128\n    sw = find_swizzle(minor_dim_bits, what='operand A')","preventionTips":["Choose tile minor dims as multiples of 128 bits (e.g. 16 x 8-bit or 8 x 16-bit elements)","Compute dtype_bits * vector_length and assert divisibility when picking block shapes","Use standard-width dtypes for MMA operands"],"tags":["jax","pallas","mosaic-gpu","swizzle","shared-memory","alignment","matmul"],"backgroundTag":"memory-alignment-not-multiple","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}