{"record":{"id":"1bdcdc8192178bcd","repo":"jax-ml/jax","slug":"compiler-params-for-platform-platform-cannot-be","errorCode":null,"errorMessage":"Compiler params for platform {platform} cannot be used for {expected_platform} lowering.","messagePattern":"Compiler params for platform (.+?) cannot be used for (.+?) lowering\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/core.py","lineNumber":186,"sourceCode":"  do nothing.\n  \"\"\"\n  return checkify.debug_check(condition, message)\n\n\n_backend_lowering_rules = {}\n\n\ndef register_lowering_rule(params_cls, rule, platform: str):\n  _backend_lowering_rules[params_cls] = (rule, platform)\n\n\ndef get_lowering_rule(params_cls, expected_platform: str):\n  rule_info = _backend_lowering_rules.get(params_cls)\n  if rule_info is None:\n    return None\n  rule, platform = rule_info\n  if platform != expected_platform:\n    raise ValueError(\n        f\"Compiler params for platform {platform} cannot be used for\"\n        f\" {expected_platform} lowering.\"\n    )\n  return rule\n\n\n@enum.unique\nclass RevisitMode(enum.Enum):\n  \"\"\"Specifies whether an output buffer supports revisiting.\n\n  By default, buffers can only be safely revisited at the next iteration\n  (immediate revisiting). If revisited at any other iteration, the buffer state\n  should be considered undefined.\n\n  If revisiting at any arbitrary iteration is required, use RevisitMode.ANY.\n  This will insert additional DMAs as needed to restore the buffer state.\n\n  Input buffers ignore revisit mode: as inputs read data from memory, their","sourceCodeStart":168,"sourceCodeEnd":204,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/core.py#L168-L204","documentation":"Pallas registers one lowering rule per compiler-params class per platform. get_lowering_rule looks up the rule for e.g. TPUCache or TritonCompilerParams and raises if the registered platform (TPU vs GPU) does not match the platform the computation is being lowered for. It guards against applying TPU kernel params to a GPU lowering and vice versa.","triggerScenarios":"Decorating a kernel with @pl.tpu_lowering(...) (or gpu_lowering/Triton) and then executing/lowering it on the other backend — e.g. writing a Triton kernel via jax.experimental.pallas.triton but running on TPU, or using TPU compiler params on GPU/CPU.","commonSituations":"Developing Pallas kernels on GPU then running on TPU (or CI on CPU); importing a pallas.triton module in a TPU job; version changes that moved decorators between modules.","solutions":["Match the lowering decorator to the backend you run on: use triton/gpu_lowering on GPU, tpu_lowering on TPU","Make the kernel import conditional on jax.default_backend()","Skip or substitute CPU reference implementations when running on unsupported backends"],"exampleFix":"# before\n@pl.tpu_lowering(compile_params)\ndef kernel(...): ...\n# run on GPU -> error\n\n# after\nfrom jax.experimental.pallas import tpu, triton\nif jax.default_backend() == 'tpu':\n    kernel = pl.pallas_call(kernel, ...).lowering(tpu.lowering(...))\nelse:\n    kernel = triton.pallas_call(kernel, ...)","handlingStrategy":"type-guard","validationCode":"import jax\nbackend = jax.default_backend()\n# choose lowering to match backend before building the kernel","typeGuard":"def pick_lowering(backend):\n    if backend == 'tpu':\n        import jax.experimental.pallas.tpu as tpu\n        return tpu.lowering\n    import jax.experimental.pallas.triton as triton\n    return triton.pallas_call","tryCatchPattern":"catch ValueError from lowering and re-dispatch to the correct backend kernel","preventionTips":["Branch kernel imports on jax.default_backend()","Run CI on the same backend as production for Pallas kernels"],"tags":["jax","pallas","backend-mismatch","gpu","tpu"],"backgroundTag":"platform-backend-mismatch","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}