{"record":{"id":"fd2b9d8d88d02a72","repo":"jax-ml/jax","slug":"jax-device-mesh-is-required-by-multimem-load-reduc","errorCode":null,"errorMessage":"JAX device mesh is required by multimem_load_reduce, but not defined.","messagePattern":"JAX device mesh is required by multimem_load_reduce, but not defined\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/primitives.py","lineNumber":5435,"sourceCode":"multimem_load_reduce_p = jax_core.Primitive(\"multimem_load_reduce\")\n\n@multimem_load_reduce_p.def_effectful_abstract_eval\ndef _multimem_load_reduce_abstract_eval(ref, *avals_flat, tree, collective_axes, reduction_op):\n  del collective_axes, reduction_op\n  _check_ref(ref, \"ref\", gpu_core.GMEM)\n  out_ref = ref\n  if tree is not None:\n    transforms = jax.tree.unflatten(tree, avals_flat)\n    out_ref = state.transform_type(transforms, ref)\n  assert isinstance(out_ref, state_types.AbstractRef)\n  return out_ref.inner_aval, {pallas_core.comms_effect}\n\n@lowering.register_lowering_rule(multimem_load_reduce_p, mgpu.LoweringSemantics.Lane)\ndef _multimem_load_reduce_lowering_rule(\n    ctx: lowering.LoweringRuleContext, ref, *transforms_leaves, tree, collective_axes, reduction_op,\n):\n  if (mesh_info := ctx.module_ctx.mesh_info) is None:\n    raise ValueError(\n        \"JAX device mesh is required by multimem_load_reduce, but not defined.\"\n    )\n  if set(collective_axes) != set(mesh_info.axis_names):\n    raise NotImplementedError(\n        \"Only collective_axes that include all JAX device mesh\"\n        f\" ({mesh_info.axis_names}) axes are supported, but got\"\n        f\" {collective_axes}\"\n    )\n  if (layout := ctx.out_layout_hint) is None:\n    raise RuntimeError(\n        \"Failed to infer the output layout of multimem_load_reduce. Please apply\"\n        \" plgpu.layout_cast to its output right after its creation.\"\n    )\n  if not isinstance(layout, (mgpu.TiledLayout, mgpu.WGStridedFragLayout)):\n    raise ValueError(\n        \"Only tiled and WG strided layouts are supported by\"\n        f\" multimem_load_reduce, but got {layout}\"\n    )","sourceCodeStart":5417,"sourceCodeEnd":5453,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/primitives.py#L5417-L5453","documentation":"The multimem_load_reduce op requires a JAX device mesh to know which devices participate in the multicast load, but the kernel was lowered with no mesh context (ctx.module_ctx.mesh_info is None). This ValueError is thrown by the Lane-semantics lowering rule before any layout work happens.","triggerScenarios":"Calling plgpu.multimem_load_reduce(...) in a pallas kernel that was compiled/launched outside of a Mesh/sharding_map context, so module_ctx.mesh_info is None.","commonSituations":"Prototyping a pallas kernel with jax.jit alone (no Mesh), then adding multimem ops; forgetting to wrap the kernel in sharding_map or map/with_mesh; running on a single device without a defined mesh.","solutions":["Wrap the kernel launch in a mesh-defining context such as sharding_map or jax's mesh utilities so mesh_info is populated","Define the mesh with the same axes you pass as collective_axes","If you do not actually need cross-device multicast, replace multimem_load_reduce with a plain load followed by an explicit reduction"],"exampleFix":"# before\nout = kernel_jit(x)  # no mesh\n# after\nwith sharding_map(kernel_jit, mesh, in_axes=..., out_axes=...):\n  out = ...  # mesh_info now available for multimem_load_reduce","handlingStrategy":"validation","validationCode":"assert mesh is not None, 'multimem ops require an active Mesh/sharding_map context'","typeGuard":null,"tryCatchPattern":"try:\n    kernel_jit(x)\nexcept ValueError as e:\n    if 'device mesh is required' in str(e):\n        with sharding_map(kernel_jit, mesh, ...):\n            ...","preventionTips":["Always launch multimem kernels inside sharding_map or a Mesh context","Add a startup assert that the mesh is defined before jitting distributed kernels"],"tags":["jax","pallas","device-mesh","multimem","missing-context"],"backgroundTag":"missing-device-mesh","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}