{"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":"When relayout_p is lowered at Warpgroup semantics on a 0-d (scalar) value, the only supported target layout is WGSplatFragLayout. Requesting any other layout for a scalar raises this ValueError.","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","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}