{"record":{"id":"059f435f961adc4b","repo":"jax-ml/jax","slug":"cannot-retrieve-the-architecture-without-an-insert","errorCode":null,"errorMessage":"Cannot retrieve the architecture without an insertion point","messagePattern":"Cannot retrieve the architecture without an insertion point","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":2436,"sourceCode":"\n  if element_bitwidth > 8:\n    return arith.muli(offset, c(element_bitwidth // 8, index_ty))\n  elif element_bitwidth < 8:\n    return arith.divsi(offset, c(8 // element_bitwidth, index_ty))\n  else:\n    return offset\n\n\n@dataclasses.dataclass(frozen=True)\nclass Arch:\n  major: int\n  minor: int\n\n\ndef get_arch() -> Arch:\n  ip = ir.InsertionPoint.current\n  if ip is None:\n    raise ValueError(\n        \"Cannot retrieve the architecture without an insertion point\"\n    )\n  block = ip.block\n  op = block.owner\n  while op is not None:\n    if op.name == \"builtin.module\":\n      arch_major = op.attributes[\"mosaic_gpu.arch_major\"]\n      arch_minor = op.attributes[\"mosaic_gpu.arch_minor\"]\n      assert isinstance(arch_major, ir.IntegerAttr)\n      assert isinstance(arch_minor, ir.IntegerAttr)\n      return Arch(arch_major.value, arch_minor.value)\n    op = op.parent\n  raise ValueError(\"Cannot retrieve the architecture: no module found\")\n\n\ndef reduce_shape(\n    shape: Sequence[int], axes: Sequence[int], keep_dims: bool = False\n) -> tuple[int, ...]:","sourceCodeStart":2418,"sourceCodeEnd":2454,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L2418-L2454","documentation":"get_arch (utils.py:2436) determines the target GPU architecture by walking up from the current MLIR insertion point to the module and reading mosaic_gpu.arch_major/arch_minor attributes. Outside of an active insertion point there is no IR context to inspect, so it refuses.","triggerScenarios":"Calling utils.get_arch() (directly or transitively, e.g. when constructing ops that depend on arch) from Python at import/tracing time with no active ir.InsertionPoint — no ongoing lowering or explicit 'with ir.InsertionPoint(block):' context.","commonSituations":"Building constants or computing shapes at module import time; calling in tests without setting up an MLIR context; helper code that runs before the mosaic lowering pipeline enters a function body.","solutions":["Move the get_arch() call inside the kernel/lowering code where an insertion point is active","Wrap the call in 'with ir.InsertionPoint(module.body): ...' after entering an ir.Context","Pass Arch explicitly to your helper instead of inferring it","Use mosaic_gpu.current_arch() or the pipeline's arch plumbing if available in your version"],"exampleFix":"# before\nARCH = utils.get_arch()  # at module scope\n# after\ndef kernel(...):\n    arch = utils.get_arch()  # inside, insertion point active\n    ...","handlingStrategy":"validation","validationCode":"assert ir.InsertionPoint.current is not None, 'call get_arch() inside a lowering'","typeGuard":null,"tryCatchPattern":"try:\n    arch = utils.get_arch()\nexcept ValueError:\n    arch = Arch(90, 0)  # fallback; must still stamp module attrs","preventionTips":["Never call arch-dependent helpers at import time","Thread Arch through kernel constructors explicitly"],"tags":["jax","mosaic-gpu","mlir","insertion-point","architecture"],"backgroundTag":"mlir-context-missing","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}