{"record":{"id":"faa0f339cadcb3b5","repo":"jax-ml/jax","slug":"gpu-launch-op-not-found","errorCode":null,"errorMessage":"gpu.launch op not found.","messagePattern":"gpu\\.launch op not found\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/dialect_lowering.py","lineNumber":2985,"sourceCode":"  return (\n      # pyrefly: ignore[missing-attribute]\n      op.OPERATION_NAME.startswith(\"mosaic_gpu.\")\n      or inference_utils.should_have_layout(op)\n      or inference_utils.should_have_transforms(op)\n      or inference_utils.should_have_tmem_layout(op)\n      # Does it have subblocks?\n      or any(bool(b) for r in op.regions for b in r)\n  )\n\n\ndef _gpu_launch_op(module: ir.Module) -> gpu.LaunchOp:\n  for op in module.body.operations:\n    for region in op.operation.regions:\n      for block in region.blocks:\n        for sub_op in block.operations:\n          if isinstance(sub_op, gpu.LaunchOp):\n            return sub_op\n  raise ValueError(\"gpu.launch op not found.\")\n\n\ndef _lowering_context(\n    module: ir.Module,\n    launch_context: lc.LaunchContext | None,\n    auto_barriers: bool,\n) -> LoweringContext:\n  \"\"\"Returns a `LoweringContext` for the given `LaunchContext`.\"\"\"\n  # TODO(bchetioui): fix tests to not have a test-only path polluting the API.\n  if launch_context is None:  # this case is used in some tests\n    return LoweringContext(None, None, None, None, None, auto_barriers, 10**9)\n\n  gpu_launch_op = _gpu_launch_op(module)\n  with ir.InsertionPoint.at_block_begin(gpu_launch_op.regions[0].blocks[0]):\n    eq = arith.CmpIPredicate.eq\n    i32 = ir.IntegerType.get_signless(32)\n    single_warp_per_block_predicate = arith.cmpi(\n        eq, utils.warp_idx(sync=False), utils.c(0, i32)","sourceCodeStart":2967,"sourceCodeEnd":3003,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/dialect_lowering.py#L2967-L3003","documentation":"The lowering pass searches the module body for the single gpu.launch op that wraps the kernel; if none exists it cannot proceed. This is an internal structure assumption of the Mosaic GPU pipeline.","triggerScenarios":"Calling the lower_mosaic_module pass on a module whose top-level function contains no gpu.launch op — e.g. hand-built IR, partially constructed modules, or a pass ordering where the launch wrapper hasn't been created yet.","commonSituations":"Running lower_mosaic_module directly on user IR without going through the normal kernel-build path; incorrect pass pipeline ordering after refactor/version changes.","solutions":["Use the standard entry point (e.g. mosaic gpu kernel launch utilities) that creates the gpu.launch wrapper before lowering","If building IR manually, wrap your func in a gpu.launch op before invoking the lowering","Check pass ordering — ensure launch materialization runs before the lowering pass"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"assert any(isinstance(o, gpu.LaunchOp) for f in module.body.operations for r in f.operation.regions for b in r.blocks for o in b.operations), 'missing gpu.launch'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Always build kernels through the standard mosaic launch utilities","Verify pass ordering before calling lower_mosaic_module on custom IR"],"tags":["mosaic","gpu","gpu-launch","pass-pipeline","mlir"],"backgroundTag":"missing-required-ir-op","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}