{"record":{"id":"0ab3337a638ab466","repo":"jax-ml/jax","slug":"no-axis-names-are-available-make-sure-you-are-usi","errorCode":null,"errorMessage":"No axis names are available. Make sure you are using `pl.core_map` with a `plgpu.Mesh`.","messagePattern":"No axis names are available\\. Make sure you are using `pl\\.core_map` with a `plgpu\\.Mesh`\\.","errorType":"exception","errorClass":"LookupError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":3383,"sourceCode":"  else:\n    raise NotImplementedError(f\"Unsupported dtype {x_aval.dtype}\")\n  kind = vector_dialect.CombiningKind.MUL\n  return _reduce_lowering_rule_wg(ctx, kind, acc, x, axes)\n\n\ndef _block_id(ctx: LoweringRuleContext, dim: gpu_dialect.Dimension) -> ir.Value:\n  result = gpu_dialect.block_id(dim)\n  cluster_size = ctx.launch_ctx.cluster_size\n  if math.prod(cluster_size) == 1 or cluster_size[dim.value] == 1:\n    return result\n  # We scale the grid in the presence of clusters, so we need to scale the\n  # block ID back here.\n  return arith_dialect.divui(result, _as_index(cluster_size[dim.value]))\n\n\ndef _resolve_cluster_axis(axis_names: _AxisNames | None, axis_name: Hashable):\n  if not axis_names:\n    raise LookupError(\n        \"No axis names are available. Make sure you are using `pl.core_map`\"\n        \" with a `plgpu.Mesh`.\"\n    )\n  if not axis_names or axis_name not in axis_names.cluster:\n    raise LookupError(\n        f\"Unknown cluster axis {axis_name}, available axes:\"\n        f\" {[*axis_names.cluster]}\"\n    )\n  return gpu_dialect.Dimension(axis_names.cluster.index(axis_name))\n\n\ndef _is_block_local_scope(collective_axes: CollectiveAxesType,\n                          axis_names: _AxisNames):\n  \"\"\"Returns whether the collective axes represents a block scope.\"\"\"\n  if axis_names.wg is None:\n    return not collective_axes\n  else:\n    return collective_axes == (axis_names.wg,)","sourceCodeStart":3365,"sourceCodeEnd":3401,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L3365-L3401","documentation":"Raised when a Pallas kernel calls a collective op with a named cluster axis but the kernel was launched without cluster axis names. Cluster axis names are only populated when the kernel is mapped with pl.core_map using a plgpu.Mesh that declares cluster axes.","triggerScenarios":"Calling e.g. mgpu.cluster_id / collective primitives with an axis_name inside a kernel launched via p.map or without plgpu.Mesh, so ctx.module_ctx.axis_names for the cluster is None/empty.","commonSituations":"Migrating kernels from single-CTA launches to multi-CTA clusters but forgetting to wrap with pl.core_map and a plgpu.Mesh; using TPU-style mesh APIs on GPU.","solutions":["Launch the kernel with pl.core_map and pass a plgpu.Mesh(axis_names=...) declaring the cluster axes","Give the mesh axis the same name you use in the collective call","If you don't need cluster collectives, remove the axis-name-based collective from the kernel"],"exampleFix":"// before\nkernel = p.map(step, in_specs=..., out_specs=...)\n// after\nmesh = plgpu.Mesh(axis_names=('row','))\nkernel = pl.core_map(step, mesh=mesh, in_specs=..., out_specs=...)","handlingStrategy":"validation","validationCode":"assert mesh is not None and mesh.axis_names, 'cluster collectives require plgpu.Mesh with axis_names'","typeGuard":null,"tryCatchPattern":"try:\n    pl.core_map(step, mesh=mesh, ...)\nexcept LookupError as e:\n    raise ValueError(f'Bad launch config: {e}') from e","preventionTips":["Always pair cluster collectives with pl.core_map + plgpu.Mesh","Keep mesh definitions next to kernel definitions","Integration-test multi-CTA launches"],"tags":["jax","pallas","mosaic-gpu","cluster","mesh","core-map"],"backgroundTag":"missing-mesh-axis-names","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}