{"record":{"id":"c725a1ab13444b8f","repo":"jax-ml/jax","slug":"expected-slice-start-start-and-slice-size-si","errorCode":null,"errorMessage":"Expected slice start ({start}) and slice size ({size}) to be divisible by the tile size ({tile})","messagePattern":"Expected slice start \\((.+?)\\) and slice size \\((.+?)\\) to be divisible by the tile size \\((.+?)\\)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":824,"sourceCode":"        indexer_shape[: -len(self.tiling)],\n        indexer_shape[-len(self.tiling) :],\n    )\n    for idx, tile, dim in zip(tiled_idxs, self.tiling, untiled_shape):\n      match idx:\n        case slice() | indexing.Slice():\n          if isinstance(idx, slice):\n            ds = indexing.Slice.from_slice(idx, dim)\n          else:\n            ds = idx\n          if ds.stride is not None and ds.stride != 1:\n            raise NotImplementedError(\n                f\"Strided slices unsupported. Got stride: {ds.stride}\"\n            )\n          start, size = ds.start, ds.size\n          if (\n              start is not None and isinstance(start, int) and start % tile\n          ) or (size is not None and isinstance(size, int) and size % tile):\n            raise ValueError(\n                f\"Expected slice start ({start}) and slice size ({size})\"\n                f\" to be divisible by the tile size ({tile})\"\n            )\n          def _maybe_cdiv_with_cast(x, y):\n            if x is None:\n              return None\n            if isinstance(x, jax.Array):\n              # If x is an int32, we need to make sure y is an int32 to avoid\n              # a dtype mismatch.\n              y = jnp.array(y, x.dtype)\n            return pallas_utils.cdiv(x, y)\n          new_start = _maybe_cdiv_with_cast(start, tile)\n          new_size = _maybe_cdiv_with_cast(size, tile)\n          idxs_after_tiling.append(indexing.Slice(new_start, new_size))\n        case _:\n          raise TypeError(f\"Unsupported index type: {type(idx)}\")\n    assert all(a % b == 0 for a, b in zip(untiled_shape, self.tiling))\n    tiled_shape = [","sourceCodeStart":806,"sourceCodeEnd":842,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L806-L842","documentation":"For a slice through a tiled dimension, TilingTransform.commute_ndindexer requires the slice's static start and size to be divisible by the tile size, so the slice can be rewritten exactly into tiled coordinates. Otherwise ValueError is raised with the offending start/size and tile.","triggerScenarios":"Slicing a tiled ref with offsets/sizes not aligned to the tile, e.g. tile size 32 and slice `ref[16:48]` (start 16 % 32 != 0) or size 48 not a multiple of 32.","commonSituations":"Block kernels with non-aligned offsets; dynamic block sizes producing unaligned slices; migrating hand-written SMEM code to tiled block specs where alignment was implicit.","solutions":["Pad or round slice start and size up/down to multiples of the tile size","Choose a tile size that divides your access pattern (e.g. tile=16 for 16-aligned blocks)","Use integer indexing on the tiled dim with computed tile indices instead of raw slices"],"exampleFix":"# before\nx = ref[16:48]  # start 16, size 32; tile=32 -> start not divisible\n\n# after\nTILE = 32\nstart = (16 // TILE) * TILE        # align to tile boundary\nx = ref[start:start + 2 * TILE]    # aligned start and tile-multiple size","handlingStrategy":"validation","validationCode":"start, size, tile = 16, 32, 32\naligned = (start % tile == 0) and (size is None or size % tile == 0)","typeGuard":"def slice_aligned_to_tile(start, size, tile) -> bool:\n    return (start is None or start % tile == 0) and (\n        size is None or size % tile == 0\n    )","tryCatchPattern":"null","preventionTips":["Design block sizes as multiples of tile size","Round offsets to tile boundaries before slicing"],"tags":["jax","pallas","alignment","tiling","slicing"],"backgroundTag":"tile-alignment-violation","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}