{"record":{"id":"1e765d2ad8b81d56","repo":"jax-ml/jax","slug":"get-global-discharge-is-not-supported-in-interpret","errorCode":null,"errorMessage":"get_global discharge is not supported in interpret mode.","messagePattern":"get_global discharge is not supported in interpret mode\\.","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/primitives.py","lineNumber":858,"sourceCode":"\n  Example::\n\n    sem_ref = pl.get_global(plgpu.SemaphoreType.REGULAR)\n    pl.semaphore_signal(sem_ref)\n    pl.semaphore_wait(sem_ref)\n  \"\"\"\n  ref_aval = what.get_ref_aval()\n  return get_global_p.bind(what=ref_aval)\n\n\n@get_global_p.def_abstract_eval\ndef _get_global_abstract_eval(*, what):\n  return what\n\n\ndef _get_global_discharge_rule(ctx, *, what):\n  del ctx, what\n  raise NotImplementedError(\n      \"get_global discharge is not supported in interpret mode.\"\n  )\n\n\nstate_discharge.register_discharge_rule(get_global_p)(\n    _get_global_discharge_rule\n)\n\n\ndef _get_ref_and_transforms(ref):\n  if isinstance(ref, state.TransformedRef):\n    return ref.ref, ref.transforms\n  return ref, ()\n\n\nclass DeviceIdType(enum.Enum):\n  MESH = \"mesh\"\n  LOGICAL = \"logical\"","sourceCodeStart":840,"sourceCodeEnd":876,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/primitives.py#L840-L876","documentation":"get_global (reading a global buffer in a Pallas kernel) has no state-discharge rule, so it cannot be used in interpret mode (or any path that requires discharging state).","triggerScenarios":"Using get_global(...) inside run_scoped or kernel code executed in interpret mode (e.g., Pallas interpretation on CPU).","commonSituations":"Debugging a TPU kernel via interpret mode that reads global buffers; running tests on CPU that use get_global.","solutions":["Avoid get_global in code paths run under interpret mode","Pass the needed values as ordinary arguments/refs instead of reading globals","Test on the actual Pallas backend rather than interpret mode"],"exampleFix":"// before\ng = get_global(...)\n// after\ndef kernel(ref, g_ref):  # pass value as a Ref argument\n    g = g_ref[...]","handlingStrategy":"type-guard","validationCode":"if INTERPRET_MODE:\n    assert not uses_get_global(f), \"get_global unsupported in interpret mode\"","typeGuard":"def kernel_is_interpret_safe(fn) -> bool:\n    src = inspect.getsource(fn)\n    return \"get_global\" not in src","tryCatchPattern":"try:\n    run_kernel(...)\nexcept NotImplementedError as e:\n    if \"get_global discharge\" in str(e):\n        skip_interpret_test()\n    raise","preventionTips":["Avoid get_global in kernels that must run in interpret mode","Pass globals as Ref arguments for testability"],"tags":["pallas","get-global","interpret-mode","not-implemented","jax"],"backgroundTag":"unsupported-feature-not-implemented","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}