{"record":{"id":"f865d90542d8c4aa","repo":"jax-ml/jax","slug":"only-interpret-mode-is-supported-on-cpu-backend","errorCode":null,"errorMessage":"Only interpret mode is supported on CPU backend.","messagePattern":"Only interpret mode is supported on CPU backend\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/pallas_call.py","lineNumber":885,"sourceCode":"      from jax._src.pallas.mosaic_gpu.interpret import params as gpu_params  # pyrefly: ignore[missing-import]\n    except ImportError:\n      pass\n    else:\n      if isinstance(interpret, gpu_params.InterpretGPUParams):\n        impl = partial(\n            mosaic_gpu_interpret.interpret_pallas_call,\n            interpret_params=interpret,\n            **params,\n        )\n\n    return mlir.lower_fun(impl, multiple_results=True)(ctx, *in_nodes)\n\n  def cpu_lowering(\n      ctx: mlir.LoweringRuleContext,\n      *in_nodes: ir.Value | Sequence[ir.Value],\n      **params,\n  ):\n    raise ValueError(\"Only interpret mode is supported on CPU backend.\")\n\n  def tpu_lowering(\n      ctx: mlir.LoweringRuleContext,\n      *in_nodes: ir.Value | Sequence[ir.Value],\n      **params,\n  ):\n    compiler_params = params.get(\"compiler_params\")\n    if compiler_params is not None:\n      rule = pallas_core.get_lowering_rule(type(compiler_params), \"tpu\")\n      if rule is not None:\n        return rule(ctx, *in_nodes, **params)\n\n    try:\n      from jax._src.pallas.mosaic import pallas_call_registration as mosaic_tpu_backend  # pyrefly: ignore[missing-import]\n    except ImportError:\n      raise _unsupported_lowering_error(\"tpu\") from None\n\n    return mosaic_tpu_backend.pallas_call_tpu_lowering_rule(","sourceCodeStart":867,"sourceCodeEnd":903,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/pallas_call.py#L867-L903","documentation":"The Pallas CPU backend only implements interpret mode — there is no real Pallas compilation pipeline for CPU. If a pallas_call is lowered on CPU without interpret=True, the registered cpu_lowering rule raises ValueError('Only interpret mode is supported on CPU Backend.') at compile time.","triggerScenarios":"Running pallas_call(..., interpret=False) (the default) on a CPU-only machine or with JAX_PLATFORMS=cpu; writing kernels for TPU/GPU and executing/testing locally on CPU.","commonSituations":"Developers without GPU/TPU hardware running Pallas examples locally; CI machines that are CPU-only; accidentally forcing the CPU backend via jax.config or JAX_PLATFORM_NAME while developing Triton/Mosaic kernels.","solutions":["Set interpret=True in pallas_call for CPU runs (debugging mode; slow but functional)","Run the code on an actual GPU/TPU backend where the kernel can be compiled","Gate hardware-specific code: use interpret=True when jax.devices()[0].platform == 'cpu'","Use interpreter mode in CI and real hardware in production"],"exampleFix":"// before\nout = pallas_call(kernel, grid, out_shape)(x, y)  # fails on CPU\n// after\nimport jax\ninterpret = jax.devices()[0].platform == 'cpu'\nout = pallas_call(kernel, grid, out_shape, interpret=interpret)(x, y)","handlingStrategy":"validation","validationCode":"import jax\nis_cpu = jax.devices()[0].platform == 'cpu'\nif is_cpu and not interpret:\n    interpret = True  # or raise early with a clear message","typeGuard":"def needs_interpret(interpret: bool) -> bool:\n    return jax.devices()[0].platform == 'cpu' and not interpret","tryCatchPattern":"try:\n    pallas_call(kernel, grid, out_shape)(x)\nexcept ValueError as e:\n    if 'interpret mode' in str(e):\n        out = pallas_call(kernel, grid, out_shape, interpret=True)(x)","preventionTips":["Parameterize interpret by device platform in all Pallas code","Use interpret=True for local CPU development and CI","Verify jax.devices() platform before compiling kernels"],"tags":["jax","pallas","cpu","backend","interpret-mode"],"backgroundTag":"kernel-backend-not-supported","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}