jax-ml/jax · error · NotImplementedError

get_global discharge is not supported in interpret mode.

Error message

get_global discharge is not supported in interpret mode.

What it means

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).

Source

Thrown at jax/_src/pallas/primitives.py:858

  Example::

    sem_ref = pl.get_global(plgpu.SemaphoreType.REGULAR)
    pl.semaphore_signal(sem_ref)
    pl.semaphore_wait(sem_ref)
  """
  ref_aval = what.get_ref_aval()
  return get_global_p.bind(what=ref_aval)


@get_global_p.def_abstract_eval
def _get_global_abstract_eval(*, what):
  return what


def _get_global_discharge_rule(ctx, *, what):
  del ctx, what
  raise NotImplementedError(
      "get_global discharge is not supported in interpret mode."
  )


state_discharge.register_discharge_rule(get_global_p)(
    _get_global_discharge_rule
)


def _get_ref_and_transforms(ref):
  if isinstance(ref, state.TransformedRef):
    return ref.ref, ref.transforms
  return ref, ()


class DeviceIdType(enum.Enum):
  MESH = "mesh"
  LOGICAL = "logical"

View on GitHub (pinned to 1e1c6a8fc0)

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

Example fix

// before
g = get_global(...)
// after
def kernel(ref, g_ref):  # pass value as a Ref argument
    g = g_ref[...]

Defensive patterns

Strategy: type-guard

Validate before calling

if INTERPRET_MODE:
    assert not uses_get_global(f), "get_global unsupported in interpret mode"

Type guard

def kernel_is_interpret_safe(fn) -> bool:
    src = inspect.getsource(fn)
    return "get_global" not in src

Try / catch

try:
    run_kernel(...)
except NotImplementedError as e:
    if "get_global discharge" in str(e):
        skip_interpret_test()
    raise

Prevention

Avoid get_global in kernels that must run in interpret mode
Pass globals as Ref arguments for testability

When it happens

Trigger: Using get_global(...) inside run_scoped or kernel code executed in interpret mode (e.g., Pallas interpretation on CPU).

Common situations: Debugging a TPU kernel via interpret mode that reads global buffers; running tests on CPU that use get_global.

Related errors

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/1e765d2ad8b81d56. Report an issue: GitHub.