{"record":{"id":"a83ddcb535959f38","repo":"jax-ml/jax","slug":"scatter-only-supports-vectorsubcoremesh-got-type","errorCode":null,"errorMessage":"Scatter only supports VectorSubcoreMesh, got {type(ref_aval.memory_space.mesh)}","messagePattern":"Scatter only supports VectorSubcoreMesh, got (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic/sc_primitives.py","lineNumber":354,"sourceCode":"          f\"{mask.shape=} does not match expected shape {expected_shape}\"\n      )\n    if mask.dtype != jnp.bool:\n      raise TypeError(f\"Mask must be a boolean array, got {mask.dtype}\")\n  effects: set[jax_core.Effect] = {state_types.WriteEffect(0)}\n  if add:\n    effects.add(state_types.ReadEffect(0))\n  return (), effects\n\n\n@sc_lowering.register_lowering_rule(scatter_p)\ndef _scatter_lowering_rule(\n    ctx: sc_lowering.LoweringRuleContext, *flat_args, tree, add\n):\n  ref, transforms, indices, x, mask = jax.tree.unflatten(tree, flat_args)\n  ref_aval, *_ = tree.unflatten(ctx.avals_in)\n  if isinstance(ref_aval.memory_space, pallas_core.CoreMemorySpace):\n    if not isinstance(ref_aval.memory_space.mesh, sc_core.VectorSubcoreMesh):\n      raise ValueError(\n          \"Scatter only supports VectorSubcoreMesh, got\"\n          f\" {type(ref_aval.memory_space.mesh)}\"\n      )\n    memory_space = ref_aval.memory_space.memory_space\n  else:\n    memory_space = ref_aval.memory_space\n  if memory_space not in (\n      tpu_core.MemorySpace.VMEM,\n      pallas_core.MemorySpace.DEFAULT,\n  ):\n    raise ValueError(\n        f\"Scatter only supports storing to VMEM, got {memory_space}\"\n    )\n  if transforms:\n    ref_block_shape, *_ = ctx.block_shapes\n    ref, _ = tc_lowering._transform_ref(\n        ref, ref_aval, ref_block_shape, transforms\n    )","sourceCodeStart":336,"sourceCodeEnd":372,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/sc_primitives.py#L336-L372","documentation":"When the scatter ref lives in a CoreMemorySpace, its mesh must be a VectorSubcoreMesh; other mesh types are unsupported by the SC scatter lowering.","triggerScenarios":"Binding scatter to a ref whose memory_space is a pallas_core.CoreMemorySpace constructed with a custom/unsupported mesh (e.g. a TensorCore mesh or new mesh type).","commonSituations":"Experimenting with new mesh APIs in Pallas; version skew where a mesh class was renamed/split; using TC-oriented memory scopes in an SC kernel.","solutions":["Construct the memory scope with sc_core.VectorSubcoreMesh (the only supported mesh)","Check the mesh type you pass to CoreMemorySpace in your kernel's memory declarations","Update/align jax version so the VectorSubcoreMesh API matches your code"],"exampleFix":"// before\nms = pallas_core.CoreMemorySpace(memory_space=..., mesh=custom_mesh)\n\n// after\nfrom jax._src.pallas.mosaic import sc_core\nms = pallas_core.CoreMemorySpace(memory_space=..., mesh=sc_core.VectorSubcoreMesh())","handlingStrategy":"validation","validationCode":"ms = ref_aval.memory_space\nif isinstance(ms, pallas_core.CoreMemorySpace):\n    assert isinstance(ms.mesh, sc_core.VectorSubcoreMesh), type(ms.mesh)","typeGuard":"def has_supported_mesh(ref_aval) -> bool:\n    ms = ref_aval.memory_space\n    return not isinstance(ms, pallas_core.CoreMemorySpace) or isinstance(ms.mesh, sc_core.VectorSubcoreMesh)","tryCatchPattern":null,"preventionTips":["Only use sc_core.VectorSubcoreMesh in SC kernels","Wrap memory scope creation in one factory function","Pin the jax version when using experimental mesh APIs"],"tags":["jax","pallas","sparsecore","mesh","scatter"],"backgroundTag":"jax-pallas-memory-space-mismatch","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}