{"record":{"id":"fc1d68d827b5a99d","repo":"jax-ml/jax","slug":"multiple-core-support-not-implemented","errorCode":null,"errorMessage":"Multiple core support not implemented.","messagePattern":"Multiple core support not implemented\\.","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/primitives.py","lineNumber":1114,"sourceCode":"    flat_device_ids = tree_util.tree_leaves(device_ids)\n    if not flat_device_ids:\n      return out\n    out = pp.concat([out, pp.text(\" \"), _pp_device_id(device_ids, context)])\n  return out\n\njax_core.pp_eqn_rules[semaphore_signal_p] = _semaphore_signal_pp_eqn\n\n\ndef _semaphore_signal_discharge_rule(ctx,\n                                     *flat_args,\n                                     args_tree,\n                                     device_id_type):\n  del device_id_type\n  [ref, transforms, inc, device_id, core_index] = args_tree.unflatten(flat_args)\n  if device_id is not None:\n    raise NotImplementedError(\"Remote signal not implemented.\")\n  if core_index is not None:\n    raise NotImplementedError(\"Multiple core support not implemented.\")\n  sem_value = _transform_semaphore(ref, transforms, ctx.in_avals[0])\n  inc = inc.astype(pallas_core.SEMAPHORE_INTERPRET_DTYPE)\n  _, new_sem_value = state_discharge.transform_swap_array(\n      ref, transforms, sem_value + inc\n  )\n  return (new_sem_value,) + (None,) * (len(ctx.in_avals) - 1), ()\nstate_discharge.register_discharge_rule(semaphore_signal_p)(\n    _semaphore_signal_discharge_rule\n)\n\n\nsemaphore_wait_p = jax_core.Primitive('semaphore_wait')\nsemaphore_wait_p.multiple_results = True\n\n\ndef semaphore_wait(\n    sem_or_view, value: int | jax_typing.Array = 1, *, decrement: bool = True\n):","sourceCodeStart":1096,"sourceCodeEnd":1132,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/primitives.py#L1096-L1132","documentation":"core_index (multi-core signaling, e.g., TPU v5p/Vegeta cores) is not supported in the discharge/interpret path of semaphore_signal.","triggerScenarios":"Passing core_index to semaphore_signal while executing under interpret mode or a discharge-requiring path.","commonSituations":"Testing multicore TPU kernel semaphore logic locally in interpret mode.","solutions":["Omit core_index in interpret-mode runs","Run multicore signaling tests only on real TPU hardware via pallas_call"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":"if INTERPRET_MODE and core_index is not None:\n    raise unittest.SkipTest(\"multicore signal unsupported in interpret mode\")","typeGuard":null,"tryCatchPattern":"try:\n    semaphore_signal(sem, 1, core_index=ci)\nexcept NotImplementedError as e:\n    if \"Multiple core\" in str(e):\n        semaphore_signal(sem, 1)\n    raise","preventionTips":["Gate multicore semaphore tests on TPU backend availability"],"tags":["pallas","semaphore","multi-core","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"}