jax-ml/jax · error · ValueError

Expected 0D shape, but got {fa.shape}

Error message

Expected 0D shape, but got {fa.shape}

What it means

For scalar-typed outputs of inline_mgpu, the returned FragmentedArray must be 0D. A non-empty shape on a value declared as scalar raises this ValueError.

Source

Thrown at jax/_src/pallas/mosaic_gpu/primitives.py:3848

    ):
      if not isinstance(fa, mgpu.FragmentedArray):
        raise ValueError(f"Expected a FragmentedArray, but got: {fa}")
      if isinstance(result_ty, ir.VectorType):
        result_shape = ir.VectorType(result_ty).shape
        if fa.shape != tuple(result_shape):
          raise ValueError(f"Expected {result_shape} but got {fa.shape}")
        if out_layout != mgpu.layouts.to_layout_attr(fa.layout):
          raise ValueError(
              f"Output layout {out_layout} does not match the layout of the"
              f" returned fragmented array {fa.layout}."
          )
        ir_ret.append(
            mgpu.dialect_lowering.fragmented_array_to_ir(fa, result_ty)
        )
      else:  # scalar case.
        assert out_layout is None
        if fa.shape:
          raise ValueError(f"Expected 0D shape, but got {fa.shape}")
        if not isinstance(fa.layout, mgpu.WGSplatFragLayout):
          raise ValueError(f"Expected WGSplatFragLayout, but got {fa.layout}")
        value = fa.registers.item()
        ir_ret.append(value)

    mgpu.dialect.return_(ir_ret)


@lowering.register_lowering_rule(inline_mgpu_p, mgpu.LoweringSemantics.Warpgroup)
@lowering.register_lowering_rule(inline_mgpu_p, *gpu_core.WGxWARP_SEMANTICS)
def _inline_mgpu_lowering_rule_wg_semantics(
    ctx: lowering.LoweringRuleContext,
    *flat_args_and_transforms,
    mgpu_fn: Callable[..., Any],
    flat_arg_types,
    flat_ret_ty,
    pytree_args,
    pytree_ref_transforms,

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Reduce the value to a scalar inside the callback, or declare the output with its true shape

Example fix

// before
ret_ty = jax.ShapeDtypeStruct((), jnp.float32)  # callback returns (4,)
// after
inline_mgpu(lambda x: jnp.sum(x), x)  # scalar returned

assert not fa.shape, fa.shape  # scalar output must be 0D

When it happens

Trigger: Declaring an output as scalar (empty-shape ShapeDtypeStruct) while the callback returns a FragmentedArray with a non-empty shape.

Common situations: Forgetting to reduce an array result to a scalar; shape bookkeeping errors in the output type list.

Related errors

• Block shape for {origin} (= {block_shape}) must have the sam
• index_map returned a value of type {type(idx_aval)} at posit
• Stacking only supported when the block size along the stack
• Sum of sizes {n} must be equal to dimension {axis} of the op
• Block size must be a multiple of the input size. Got block {

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/8d6fd06c8a70fbec. Report an issue: GitHub.