{"record":{"id":"45ee2cb3facea1c7","repo":"jax-ml/jax","slug":"non-indexing-transforms-on-gmem-refs-are-not-imple","errorCode":null,"errorMessage":"Non-indexing transforms on GMEM refs are not implemented.","messagePattern":"Non-indexing transforms on GMEM refs are not implemented\\.","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/primitives.py","lineNumber":460,"sourceCode":"            \"Multicast refs are not supported by this primitive.\"\n        )\n      if (mesh_info := ctx.module_ctx.mesh_info) is None:\n        raise ValueError(\n            \"JAX device mesh is required by multicast copies, but not defined.\"\n            \" Use jax.set_mesh.\"\n        )\n      if set(transform.collective_axes) != set(mesh_info.axis_names):\n        raise NotImplementedError(\n            \"Only collective_axes that include all JAX device mesh  axes are\"\n            f\" supported, but got {transform.collective_axes}. Make sure to\"\n            f\" pass collective_axes={mesh_info.axis_names}\"\n        )\n      peer_id = mgpu.GLOBAL_BROADCAST\n      continue\n    elif isinstance(transform, indexing.NDIndexer):\n      indexers.append(transform)\n    else:\n      raise NotImplementedError(\n          \"Non-indexing transforms on GMEM refs are not implemented.\")\n  if indexers:\n    indexer = lowering.merge_indexers(indexers)\n    gmem_slice = lowering._ndindexer_indices(indexer, allow_arrays=True)\n  else:\n    gmem_slice = ()\n  return dict(\n      gmem_slice=gmem_slice,\n      gmem_peer_id=peer_id,\n  )\n\n\ndef _extract_smem_copy_params(aval, transforms):\n  if not transforms:\n    return {}\n  # Split off swizzling, if present\n  match transforms:\n    case [gpu_core.UnswizzleRef(swizzle), *transforms]:","sourceCodeStart":442,"sourceCodeEnd":478,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/primitives.py#L442-L478","documentation":"While walking transforms on a GMEM reference, _extract_gmem_copy_params only understands NDIndexer transforms (plus the multicast transform handled above). Any other transform type on the ref (slicing helpers, swizzle, permutation, etc.) reaches the else branch and raises NotImplementedError.","triggerScenarios":"Passing a GMEM BlockRef with a non-indexing transform (e.g. a sliced/swizzled/permuted ref produced by ref.swap_dims, unswizzle, or similar) to copy_gmem_to_smem/copy_smem_to_gmem/prefetch in a Mosaic GPU kernel.","commonSituations":"Trying to reuse SMEM-style layout transforms (tiled/swizzled refs) on GMEM refs; applying fancy indexing utilities not based on indexing.NDIndexer; upgrading JAX where transform kinds expanded beyond what the copy path handles.","solutions":["Index GMEM refs only with standard Pallas indexing (block_indices / NDIndexer-producing ops) before the copy","Materialize the transformed value into an intermediate SMEM buffer and copy that instead","Check the JAX version/changelog — support for additional transforms may have been added in newer releases"],"exampleFix":"# before\ncopy_gmem_to_smem(gref.unswizzle(swizzle), smem)\n# after\nsmem[...] = gref[...]  # plain indexing\ncopy_gmem_to_smem(gref, smem)","handlingStrategy":"type-guard","validationCode":null,"typeGuard":"from jax._src.pallas.mosaic_gpu import indexing\ndef refs_have_only_indexer_transforms(ref_transforms):\n  return all(t is None or isinstance(t, indexing.NDIndexer) or _is_multicast(t)\n             for t in ref_transforms)","tryCatchPattern":null,"preventionTips":["Keep GMEM refs limited to plain NDIndexer indexing before copies","Materialize layout-transformed data into SMEM before GMEM copies"],"tags":["jax","pallas","mosaic-gpu","gmem-ref","transforms","not-implemented"],"backgroundTag":"unsupported-operation-argument","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}