{"record":{"id":"d827c2ab97a22226","repo":"jax-ml/jax","slug":"non-decrementing-wait-is-not-supported","errorCode":null,"errorMessage":"Non-decrementing wait is not supported.","messagePattern":"Non-decrementing wait is not supported\\.","errorType":"validation","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic/interpret/interpret_pallas_call.py","lineNumber":1736,"sourceCode":"            ctx.axis_indices)\n        token = callback.io_callback(\n            functools.partial(semaphore_signal, source_info=eqn.source_info),\n            TOKEN_SHAPE_DTYPE,\n            token,\n            ctx.device_id,\n            ctx.local_core_id,\n            state_discharge.transform_array(sem, sem_transforms),\n            inc,\n            target_device_id,\n            core_index,\n        )\n        out = []\n\n      elif prim is primitives.semaphore_wait_p:\n        sem, sem_transforms, value, decrement = (\n            jax.tree.unflatten(eqn.params['args_tree'], deferred_invals()))\n        if not decrement:\n          raise NotImplementedError('Non-decrementing wait is not supported.')\n        token = callback.io_callback(\n            semaphore_wait,\n            TOKEN_SHAPE_DTYPE,\n            token,\n            ctx.device_id,\n            ctx.local_core_id,\n            state_discharge.transform_array(sem, sem_transforms),\n            value,\n        )\n        out = []\n\n      else:\n        if ctx.interpret_params.skip_floating_point_ops and all(\n            interpret_utils.is_float(ovar.aval.dtype) for ovar in eqn.outvars\n        ):\n          # Skip `prim.bind` since `prim` only produces floating-point values.\n          # It is safe to populate `out` with avals since mapping `write` over\n          #  `out` below only relies on the shape and dtype (for writing","sourceCodeStart":1718,"sourceCodeEnd":1754,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic/interpret/interpret_pallas_call.py#L1718-L1754","documentation":"The Pallas TPU interpreter only supports semaphore waits that decrement the semaphore's value; semaphore_wait with decrement=False (a non-destructive read of the semaphore) is not implemented and raises NotImplementedError.","triggerScenarios":"Calling primitives.semaphore_wait / pltpu.semaphore_wait with decrement=False inside a kernel run in TPU interpret mode.","commonSituations":"Polling a semaphore without consuming it (common in spin-wait patterns); GPU-style atomic semantics ported to TPU; probing semaphore state for debugging.","solutions":["Use decrementing waits (decrement=True), ensuring each wait pairs with a signal/lock post","Restructure spin-wait patterns to signal/wait discipline instead of non-destructive polling","Fall back to real TPU execution if non-decrementing waits are required"],"exampleFix":"# before\npltpu.semaphore_wait(sem, value=0, decrement=False)\n# after\npltpu.semaphore_wait(sem, value=0)  # decrement=True default","handlingStrategy":"fallback","validationCode":"assert decrement is not False, 'non-decrementing semaphore wait unsupported in TPU interpret mode'","typeGuard":"null","tryCatchPattern":"try:\n    interpret_run(kernel)\nexcept NotImplementedError as e:\n    if 'Non-decrementing wait' in str(e):\n        run_on_tpu_hardware(kernel)","preventionTips":["Use decrementing waits paired with signals","Avoid polling semaphore values without decrement","Test semaphore patterns in interpret mode before scaling"],"tags":["jax","pallas","tpu","semaphore","interpret-mode","not-implemented"],"backgroundTag":"unsupported-feature","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}