{"record":{"id":"4368e1e423500c38","repo":"jax-ml/jax","slug":"get-global-only-supports-semaphores-got-what","errorCode":null,"errorMessage":"get_global only supports semaphores, got {what}","messagePattern":"get_global only supports semaphores, got (.+?)","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/lowering.py","lineNumber":3716,"sourceCode":"          jaxpr,\n          input_refs,\n          consts,\n      )\n\n  assert len(outs) == len(jaxpr.outvars), (jaxpr, outs)\n  return outs\n\n\n@_register_resource_estimator(primitives.get_global_p)\ndef _get_global_resource_estimator(\n    ctx: ResourceEstimatorContext, *, what\n) -> Resources:\n  if what.memory_space == gpu_core.GMEM and jnp.issubdtype(\n      what.dtype, pallas_core.semaphore\n  ):\n    collective_axes = tuple(ctx.axis_names)\n    return Resources(scoped_gmem_semaphores={collective_axes: what.size})\n  raise NotImplementedError(f\"get_global only supports semaphores, got {what}\")\n\n\n@register_lowering_rule(primitives.get_global_p, mgpu.LoweringSemantics.Lane)\n@register_lowering_rule(\n    primitives.get_global_p, mgpu.LoweringSemantics.Warpgroup\n)\ndef _get_global_lowering_rule(ctx: LoweringRuleContext, *, what):\n  if what.memory_space == gpu_core.GMEM and jnp.issubdtype(\n      what.dtype, pallas_core.semaphore\n  ):\n    collective_axes = tuple(ctx.module_ctx.axis_names)\n    return ctx.module_ctx.reserve_semaphores(\n        what.shape, collective_axes=collective_axes\n    ).__enter__()\n  raise NotImplementedError(f\"get_global only supports semaphores, got {what}\")\n\n\n@register_lowering_rule(discharge.run_state_p, mgpu.LoweringSemantics.Lane)","sourceCodeStart":3698,"sourceCodeEnd":3734,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/lowering.py#L3698-L3734","documentation":"pl.get_global (resource computation) only works for GMEM-resident semaphores: it must have memory_space GMEM and dtype pallas semaphore. Requesting anything else (e.g. an SMEM value or a regular array) raises NotImplementedError when computing kernel resources.","triggerScenarios":"Calling pl.get_global on a global whose dtype is not a semaphore or whose memory space is SMEM, during resource analysis of a GPU kernel.","commonSituations":"Trying to read arbitrary global arrays via get_global instead of block arguments; declaring a semaphore in the wrong memory space in the kernel signature.","solutions":["Only use pl.get_global for semaphores allocated in GMEM","Pass regular global arrays as normal BlockSpec'd inputs instead","Check value.memory_space == gpu_core.GMEM and jnp.issubdtype(dtype, pallas_core.semaphore) before calling"],"exampleFix":"// before\nx = pl.get_global(ordinary_array_global)\n// after\nx = kernel_input_arg  # declared via in_specs\n# get_global reserved for:\n# sem = pl.get_global(semaphore_in_gmem)","handlingStrategy":"type-guard","validationCode":"assert what.memory_space == gpu_core.GMEM and jnp.issubdtype(what.dtype, pallas_core.semaphore)","typeGuard":"def is_gmem_semaphore(aval) -> bool:\n    return (aval.memory_space == gpu_core.GMEM\n            and jnp.issubdtype(aval.dtype, pallas_core.semaphore))","tryCatchPattern":null,"preventionTips":["Use get_global only for semaphores","Route arrays through kernel inputs"],"tags":["jax","pallas","get-global","semaphore"],"backgroundTag":"unsupported-memory-space","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}