{"record":{"id":"792f645452f2f9fc","repo":"jax-ml/jax","slug":"unsupported-barrier-type-type-barrier","errorCode":null,"errorMessage":"Unsupported barrier type: {type(barrier)}","messagePattern":"Unsupported barrier type: (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py","lineNumber":942,"sourceCode":"\n  barrier, clock = shared_memory.get_barrier_and_increment_clock(\n      barrier_key, thread\n  )\n  if isinstance(barrier, memory.ClusterBarrier):\n    barrier.arrive(\n        mesh_location=mesh_location,\n        thread=thread,\n        clock=clock,\n        logging_info=memory.GPULoggingInfo(mesh_location, thread, source_info),\n    )\n  elif isinstance(barrier, memory.Barrier):\n    barrier.arrive(\n        thread,\n        clock,\n        memory.GPULoggingInfo(mesh_location, thread, source_info),\n    )\n  else:\n    raise ValueError(f\"Unsupported barrier type: {type(barrier)}\")\n  return token\n\n\n# Note that this callback is also used for arriving at cluster barriers.\ndef call_barrier_arrive(\n    token: jax.Array,\n    mesh_location: memory.MeshLocation,\n    thread: memory.Thread,\n    allocation_key_as_array: jax.Array,\n    source_info: source_info_util.SourceInfo | None = None,\n):\n  return callback.io_callback(\n      functools.partial(_barrier_arrive, source_info=source_info),\n      TOKEN_SHAPE_DTYPE,\n      token,\n      mesh_location,\n      thread,\n      allocation_key_as_array,","sourceCodeStart":924,"sourceCodeEnd":960,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/pallas/mosaic_gpu/interpret/gpu_callbacks.py#L924-L960","documentation":"The interpreter's barrier-arrive callback only supports the standard memory.Barrier implementation. Arriving at any other synchronization primitive type (e.g. a named barrier with different semantics) is rejected.","triggerScenarios":"Calling barrier_arrive / arriving at a barrier object in interpret mode whose runtime type isn't memory.Barrier (unsupported barrier variant such as cgbarrier or future barrier kinds).","commonSituations":"Using experimental barrier APIs newly added to Mosaic GPU before interpret-mode support catches up; version mismatch between kernel code and jax interpreter.","solutions":["Restructure the kernel to use plain Barrier objects for arrive semantics","Update jax / mosaic plugin so the barrier type is supported in interpret mode","Run on device instead of interpret mode while using the exotic barrier","Report upstream if the barrier type should be supported"],"exampleFix":null,"handlingStrategy":"type-guard","validationCode":null,"typeGuard":"def is_supported_barrier(b) -> bool:\n    from jax._src.pallas.mosaic_gpu.interpret import memory\n    return isinstance(b, memory.Barrier)","tryCatchPattern":"try:\n    kernel(x)\nexcept ValueError as e:\n    if 'Unsupported barrier type' in str(e):\n        swap to a plain Barrier","preventionTips":["Use only standard Barrier objects for interpret-mode testing"],"tags":["pallas","mosaic-gpu","interpret-mode","barrier","unsupported"],"backgroundTag":"unsupported-operation-in-interpreter","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}