{"record":{"id":"a031904ef7a8ded1","repo":"jax-ml/jax","slug":"inline-mgpu-p-return-type-tree-mismatch-ret","errorCode":null,"errorMessage":"inline_mgpu_p return type tree mismatch: {ret} != {return_type}","messagePattern":"inline_mgpu_p return type tree mismatch: (.+?) != (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/primitives.py","lineNumber":3671,"sourceCode":"            f\" types. Got shape={r.shape}.\"\n        )\n\n  flat_transformed = _inline_mgpu_flat_transformed_args(\n      ctx,\n      flat_args_and_transforms,\n      flat_arg_types,\n      pytree_args,\n      pytree_ref_transforms,\n  )\n  args = jax.tree.unflatten(pytree_args, flat_transformed)\n  ret = mgpu_fn(ctx.launch_ctx, *args)\n  ret_leaves, ret_tree = jax.tree.flatten(\n      ret, lambda x: isinstance(x, mgpu.FragmentedArray)\n  )\n\n  if ret_tree != pytree_ret_ty:\n    return_type = jax.tree.unflatten(pytree_ret_ty, flat_ret_ty)\n    raise ValueError(\n        f\"inline_mgpu_p return type tree mismatch: {ret} != {return_type}\"\n    )\n\n  for ty, r in zip(flat_ret_ty, ret_leaves):\n    _type_check_mgpu_lane_semantics(r, ty)\n\n  return ret_leaves\n\n\ndef _ref_type_to_transforms(ref_type: RefType) -> ir.ArrayAttr:\n  \"\"\"Returns the Mosaic GPU transforms for the given ref type.\"\"\"\n  transform_attrs = [gpu_core.to_transform_attr(t)\n                     for t in ref_type.transforms]\n  return ir.ArrayAttr.get(transform_attrs)\n\n\ndef _custom_primitive_in_specs(\n    ctx: lowering.LoweringRuleContext,","sourceCodeStart":3653,"sourceCodeEnd":3689,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/primitives.py#L3653-L3689","documentation":"The pytree structure of values returned by the user's inline_mgpu function must exactly match the declared return type tree. Flattening the actual return with FragmentedArray leaves yields a different tree structure than pytree_ret_ty, so this ValueError is raised.","triggerScenarios":"The callback returns a differently nested structure (e.g. a tuple vs a single array, wrong leaf count) than the avals passed as flat_ret_ty / their pytree structure.","commonSituations":"Returning (array,) vs array; changing the callback signature or return shape without updating the declared output avals; conditional returns producing inconsistent trees.","solutions":["Make the callback return exactly the declared pytree structure (same nesting and leaf count)","Update the declared return-type tree whenever the callback's return expression changes","Avoid conditional returns with different structures; return a fixed tuple shape"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":"import jax\nassert jax.tree.structure(fn(*args)).num_leaves == len(flat_ret_ty)","typeGuard":"def ret_tree_matches(fn, args, ret_tree):\n    return jax.tree.structure(fn(*args)) == ret_tree","tryCatchPattern":null,"preventionTips":["Return a fixed-structure tuple from callbacks","Update declared out types whenever the callback return changes"],"tags":["jax","pallas","pytree","return-type-mismatch"],"backgroundTag":"pytree-structure-mismatch","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}