{"record":{"id":"939a40a2e43ed79e","repo":"jax-ml/jax","slug":"expected-an-ndindexer-but-got-transforms-0","errorCode":null,"errorMessage":"Expected an `NDIndexer`, but got {transforms[0]}","messagePattern":"Expected an `NDIndexer`, but got (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/jaxpr_interpret.py","lineNumber":135,"sourceCode":"    transforms_treedef, transforms_leaves,\n) -> indexing.DimIndexer | None:\n  # TODO(nrink): The working out of `transforms` and the returned index below\n  # may need tidying up. Specifically, GPU interpret mode should correctly\n  # support legal ways to index into barriers. (Here, 'legal' is to be read as\n  # 'allowed by the Pallas GPU semantics'.)\n  if transforms_treedef is None:\n    return None\n  transforms = jax.tree.unflatten(transforms_treedef, transforms_leaves)\n\n  if not transforms:\n    return None\n  if not hasattr(transforms, \"__len__\") or len(transforms) != 1:\n    raise NotImplementedError(\n        f\"Indexing barrier with {transforms} not supported in GPU interpret\"\n        \" mode\"\n    )\n  if not isinstance(transforms[0], indexing.NDIndexer):\n    raise ValueError(f\"Expected an `NDIndexer`, but got {transforms[0]}\")\n  if len(transforms[0].indices) == 1:\n    return transforms[0].indices[0]\n  return tuple(transforms[0].indices)\n\n\ndef _get_barrier_allocation_key_from_inval(\n    inval, transforms_treedef, transforms_leaves\n) -> jax.Array:\n  # `inval` is expected to correspond to a barrier. Since we are interpreting,\n  # `inval` will in fact contain the allocation key (which is a Jax array) for\n  # the barrier.\n  allocation_key_as_array = inval\n\n  # Assert to check internal consistency: `allocation_key_as_array` should be\n  # at least a 2D array, and the size of the last dimension is 5 (which matches the\n  # fields count of HostAllocationKey).\n  assert len(allocation_key_as_array.shape) >= 2\n  assert (","sourceCodeStart":117,"sourceCodeEnd":153,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/jaxpr_interpret.py#L117-L153","documentation":"Barrier indexing in interpret mode only accepts transforms that are indexing.NDIndexer instances. Any other transform type on a barrier reference is rejected with this ValueError.","triggerScenarios":"Applying a non-NDIndexer transform (e.g. UnswizzleRef, UntilingTransform, or any custom Transform) to a barrier before barrier_arrive/wait, copy_gmem_to_smem, or tcgen05 MMA ops in interpret mode.","commonSituations":"Copying transform pipelines written for SMEM buffers onto barrier refs; generic code that applies unswizzle transforms to all refs in a kernel.","solutions":["Remove non-indexer transforms from barrier references","Apply unswizzle/untile transforms only to regular SMEM buffers, never to ClusterBarrierType refs","Verify the ref passed to barrier ops is the raw allocation"],"exampleFix":"// before\nop(unswizzle(barrier_ref))\n// after\nop(barrier_ref)","handlingStrategy":"type-guard","validationCode":null,"typeGuard":"from jax._src.pallas import indexing\ndef barrier_transforms_are_valid(transforms) -> bool:\n    return all(isinstance(t, indexing.NDIndexer) for t in transforms)","tryCatchPattern":null,"preventionTips":["Never apply unswizzle/untile to ClusterBarrierType refs","Keep transform pipelines for buffers separate from barriers"],"tags":["jax","mosaic-gpu","barrier","type-validation"],"backgroundTag":"invalid-argument-type","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}