{"record":{"id":"cc90a146dd9cca79","repo":"jax-ml/jax","slug":"a-custom-return-op-must-terminate-the-block","errorCode":null,"errorMessage":"A custom return op must terminate the block.","messagePattern":"A custom return op must terminate the block\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/dialect_lowering.py","lineNumber":2636,"sourceCode":") -> Sequence[ir.Value]:\n  \"\"\"Lowering rule for mgpu.CustomPrimitiveOp.\"\"\"\n  del ctx\n  block = op.body.blocks[0]\n  for arg, operand in zip(block.arguments, op.operands, strict=True):\n    arg.replace_all_uses_with(operand)\n\n  return_op = None\n  ip = ir.InsertionPoint.current\n  for block_op in block.operations:\n    if isinstance(block_op.opview, mgpu.ReturnOp):\n      assert return_op is None\n      return_op = block_op.opview\n      continue\n    block_op.detach_from_parent()\n    ip.insert(block_op)\n\n  if return_op is None:\n    raise ValueError(\"A custom return op must terminate the block.\")\n\n  return return_op.operands\n\n\n@_register_lowering(mgpu.WarpMapOp)\ndef _mgpu_warp_map_op_lowering_rule(\n    ctx: LoweringContext, op: mgpu.WarpMapOp\n) -> Sequence[ir.Value]:\n  \"\"\"Lowering rule for mgpu.WarpMapOp.\"\"\"\n  for a, o in zip(op.body.arguments, op.operands, strict=True):\n    a.replace_all_uses_with(o)\n  warp_ctx = dataclasses.replace(ctx, thread_semantics=utils.ThreadSubset.WARP)\n  # We allow the warps to schedule async copies without synchronizing with\n  # other warps, so we need to add a barrier here to make sure all reads and\n  # writes have completed.\n  if ctx.auto_barriers:\n    utils.warpgroup_barrier()\n  ip = ir.InsertionPoint.current","sourceCodeStart":2618,"sourceCodeEnd":2654,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/dialect_lowering.py#L2618-L2654","documentation":"When lowering an mgpu custom primitive, Mosaic relocates the ops from the primitive's body region into the caller and expects the region to end with a custom return op. If none was found, the body is malformed.","triggerScenarios":"Defining a custom primitive (warp-level custom op) whose body block does not terminate with mgpu.return / the custom return op — e.g. empty body or body ending in another op.","commonSituations":"Hand-written custom primitive builders; bugs generating the body IR; forgetting the mandatory terminator when constructing the region programmatically.","solutions":["Ensure the custom primitive body block ends with the required return op carrying the outputs","Use the provided helper/builder APIs that insert the terminator automatically","Validate the region has exactly one terminator before lowering"],"exampleFix":"// before\nbody = build_body_without_terminator(primitive)\n// after\nbody = build_body(primitive)\nmgpu.return_(results)  # terminator required at end of body","handlingStrategy":"validation","validationCode":"from jax.experimental.mosaic.gpu import mgpu\nterminators = [o for o in body.ops if isinstance(o.opview, mgpu.ReturnOp)]\nassert terminators, 'custom primitive body must end with return'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Always use builder helpers that insert the return terminator","Validate region structure (single terminator) before registering a custom primitive"],"tags":["mosaic","gpu","custom-primitive","mlir","terminator"],"backgroundTag":"missing-block-terminator","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}