{"record":{"id":"be9cf4f5e607cc84","repo":"jax-ml/jax","slug":"strided-slices-unsupported-got-stride-ds-stride","errorCode":null,"errorMessage":"Strided slices unsupported. Got stride: {ds.stride}","messagePattern":"Strided slices unsupported\\. Got stride: (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/core.py","lineNumber":817,"sourceCode":"    del aval\n    idxs = indexer.indices\n    indexer_shape = indexer.shape\n    untiled_idxs = idxs[: -len(self.tiling)]\n    tiled_idxs = idxs[-len(self.tiling) :]\n    idxs_after_tiling: list[indexing.Slice] = []\n    leading_shape, untiled_shape = (\n        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)","sourceCodeStart":799,"sourceCodeEnd":835,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/core.py#L799-L835","documentation":"TilingTransform.commute_ndindexer rewrites an indexer through a tiled view. Slicing is only supported with unit stride; a slice with stride None (dynamic) or any value != 1 raises NotImplementedError because strided access through tiled layouts is not implemented.","triggerScenarios":"Indexing a tiled ref with a strided slice such as `ref[::2]` or `ref[0:128:4]`, or an indexing.Slice with dynamic stride, when the tiling transform must commute the indexer during lowering.","commonSituations":"Downsampling/strided access patterns in Pallas GPU kernels; converting existing indexing code to tiled block specs; passing slices with symbolic strides that default to non-1.","solutions":["Replace the strided slice with explicit indices (e.g. `ref[jnp.arange(0, n, 2)]` via arange indexing) if supported by the surrounding op","Materialize the strided read outside the tiled ref (copy to a normal array first)","Ensure slices have stride exactly 1 (or None-with-default-1 after fixing slice construction)"],"exampleFix":"# before\nx = ref[:, ::2]  # NotImplementedError: strided slice through tiling\n\n# after\nidx = jnp.arange(0, ref.shape[1], 2)\nx = ref[:, idx]","handlingStrategy":"fallback","validationCode":"def slice_is_unit_stride(s):\n    return getattr(s, 'stride', None) in (None, 1) or getattr(s, 'step', None) in (None, 1)","typeGuard":"null","tryCatchPattern":"null","preventionTips":["Avoid ::k slicing on tiled refs","Use arange-based integer indexing for strided access"],"tags":["jax","pallas","slicing","strides","tiling"],"backgroundTag":"unsupported-strided-slice","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}