{"record":{"id":"9065376774753867","repo":"jax-ml/jax","slug":"only-single-thread-arrival-is-supported-on-pre-hop","errorCode":null,"errorMessage":"Only single-thread arrival is supported on pre-Hopper hardware","messagePattern":"Only single-thread arrival is supported on pre-Hopper hardware","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":1205,"sourceCode":"        # We need to synchronize the threads after `::before_thread_sync`, as\n        # not all threads arrive on the barrier.\n        if scope == ThreadSubset.WARPGROUP:\n          warpgroup_barrier()\n        elif scope == ThreadSubset.WARP:\n          warp_barrier()\n        else:\n          raise ValueError(f\"Unsupported scope: {scope}\")\n\n    ptx_scope = self._ptx_scope\n    if can_complete or ptx_scope != \"cta\":\n      pred_ptx = pred_constraint = \"\"\n      if predicate is not None:\n        pred_ptx = \"@$2\"\n        pred_constraint = \",b\"\n      count_ptx = f\", {arrival_count}\"\n      if get_arch().major < 9:\n        if arrival_count != 1:\n          raise ValueError(\n              \"Only single-thread arrival is supported on pre-Hopper hardware\"\n          )\n        count_ptx = \"\"\n      llvm.inline_asm(\n          ir.IntegerType.get_signless(64),\n          [self.get_ptr()] + ([predicate] if predicate is not None else []),\n          f\"{pred_ptx} mbarrier.arrive.release.{ptx_scope}.shared::{ptx_scope}.b64 $0, [$1]{count_ptx};\",\n          \"=l,r\" + pred_constraint,\n          has_side_effects=True,\n      )\n    else:\n      if predicate is not None:\n        raise NotImplementedError(\n            \"Predicate not supported for no-complete arrive\"\n        )\n      count = c(arrival_count, ir.IntegerType.get_signless(32))\n      nvvm.mbarrier_arrive_nocomplete(self.get_ptr(), count)\n","sourceCodeStart":1187,"sourceCodeEnd":1223,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L1187-L1223","documentation":"On architectures before Hopper (sm < 90), the inline PTX mbarrier.arrive path cannot carry an arrival_count operand, so only arrival_count == 1 is emitted. Requesting multi-thread arrival on pre-Hopper GPUs raises ValueError.","triggerScenarios":"barrier.arrive(arrival_count=4) on Ampere or earlier (get_arch().major < 9); porting a Hopper pipeline (e.g. TMA warp-specialized code) to A100/sm80.","commonSituations":"Developing on H100 and running on older GPUs in CI; defaulting arrival_count to warp size for producer arrival; hardcoding multi-thread arrivals copied from Blackwell examples.","solutions":["Make arrival_count configurable and set 1 on pre-Hopper (each thread arrives individually)","Gate the pipeline on jax.experimental.mosaic.gpu.utils.get_arch().major >= 9 and fall back to a single-producer scheme","Have one lane arrive on behalf of the group using a predicate on pre-Hopper hardware"],"exampleFix":"# before\nbarrier.arrive(arrival_count=threads_per_cta)\n# after\nfrom jax.experimental.mosaic.gpu.utils import get_arch\ncount = threads_per_cta if get_arch().major >= 9 else 1\nbarrier.arrive(arrival_count=count)","handlingStrategy":"validation","validationCode":"from jax.experimental.mosaic.gpu.utils import get_arch\nassert get_arch().major >= 9 or arrival_count == 1","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Make arrival_count arch-dependent: 1 on pre-Hopper, warp size on Hopper+","Test pipelines on the oldest GPU in your fleet"],"tags":["jax","mosaic-gpu","barrier","gpu-architecture","hopper"],"backgroundTag":"gpu-arch-feature-unsupported","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}