{"record":{"id":"c74c8eb34401bdc9","repo":"jax-ml/jax","slug":"only-plgpu-layout-wg-splat-is-supported-for-scalar","errorCode":null,"errorMessage":"Only plgpu.Layout.WG_SPLAT is supported for scalar values.","messagePattern":"Only plgpu\\.Layout\\.WG_SPLAT is supported for scalar values\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":4654,"sourceCode":"      \"relayout_p is not supported with Lane semantics.\"\n  )\n\n\n@register_lowering_rule(pjit.relayout_p, mgpu.LoweringSemantics.Warpgroup)\ndef _relayout_lowering_wg(\n    ctx: LoweringRuleContext, x, *, dst_layout\n):\n  if dst_layout is jax_layout.AutoLayout:\n    return x\n  layout = fa.TiledLayout(\n      dst_layout.tiling,\n      dst_layout.warp_dims,\n      dst_layout.lane_dims,\n      dst_layout.vector_dim,\n  )\n  if ctx.avals_in[0].ndim == 0:  # scalar case\n    if layout != mgpu.WGSplatFragLayout():\n      raise ValueError(\n          \"Only plgpu.Layout.WG_SPLAT is supported for scalar values.\"\n      )\n    return x\n  return mgpu.dialect.layout_cast(x, mgpu.to_layout_attr(layout))\n\n\n@register_lowering_rule(gpu_core.layout_cast_p, mgpu.LoweringSemantics.Lane)\ndef _layout_cast_lowering(ctx: LoweringRuleContext, x, *, new_layout):\n  del ctx  # Unused.\n  return x.to_layout(new_layout.to_mgpu())\n\n\n@register_lowering_rule(gpu_core.layout_cast_p, mgpu.LoweringSemantics.Warpgroup)\ndef _layout_cast_lowering_wg(\n    ctx: LoweringRuleContext, x, *, new_layout\n):\n  layout = new_layout.to_mgpu()\n  if ctx.avals_in[0].ndim == 0:  # scalar case","sourceCodeStart":4636,"sourceCodeEnd":4672,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L4636-L4672","documentation":"In the Warpgroup-semantics lowering of relayout_p, a 0-ndim (scalar) input has no tiling, warp, or lane structure, so the only faithfully representable layout is a warpgroup-splatted fragment (WGSplatFragLayout). Requesting any other plgpu.Layout as dst_layout for a scalar triggers this error.","triggerScenarios":"Calling lax.layout_cast / relayout with a non-WG_SPLAT layout on a scalar inside a Warpgroup-semantic Mosaic GPU kernel.","commonSituations":"Generic code that applies a fixed tiling layout to all values, including scalars like loop counters or accumulated scalars.","solutions":["Skip layout_cast for scalars (only cast arrays with ndim>0)","Use plgpu.Layout.WG_SPLAT for scalar values"],"exampleFix":"# before\nx = plgpu.layout_cast(scalar_val, tiled_layout)\n# after\nx = scalar_val  # scalars need no cast; or use WG_SPLAT","handlingStrategy":"type-guard","validationCode":"if x.ndim == 0: assert layout == plgpu.Layout.WG_SPLAT","typeGuard":"def needs_layout_cast(x) -> bool:\n    return getattr(x, 'ndim', 0) > 0","tryCatchPattern":null,"preventionTips":["Only cast non-scalars; use WG_SPLAT for scalars"],"tags":["pallas","mosaic-gpu","layout","scalar"],"backgroundTag":"invalid-layout","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","contentChangedAt":null,"schemaVersion":2},"datasetVersion":"2026-09-23T08:17:48.524Z"}