{"record":{"id":"a1c8f25c4e75a460","repo":"jax-ml/jax","slug":"arrive-expect-tx-is-only-supported-on-hopper-hard","errorCode":null,"errorMessage":"arrive_expect_tx is only supported on Hopper+ hardware","messagePattern":"arrive_expect_tx is only supported on Hopper\\+ hardware","errorType":"validation","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":1228,"sourceCode":"          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\n  def arrive_expect_tx(\n      self, bytes: int | ir.Value, predicate: ir.Value | None = None\n  ):\n    if get_arch().major < 9:\n      raise NotImplementedError(\"arrive_expect_tx is only supported on Hopper+ hardware\")\n    if isinstance(bytes, int):\n      bytes = c(bytes, ir.IntegerType.get_signless(32))\n    elif isinstance(bytes.type, ir.IndexType):\n      i32 = ir.IntegerType.get_signless(32)\n      bytes = arith.index_cast(i32, bytes)\n    nvvm.mbarrier_arrive_expect_tx(\n        self.get_ptr(), bytes, predicate=predicate, scope=self._nvvm_scope\n    )\n\n  def complete_tx(\n      self, bytes: int | ir.Value, predicate: ir.Value | None = None\n  ):\n    if get_arch().major < 9:\n      raise NotImplementedError(\"complete_tx is only supported on Hopper+ hardware\")\n    if isinstance(bytes, int):\n      bytes = c(bytes, ir.IntegerType.get_signless(32))\n    elif isinstance(bytes.type, ir.IndexType):\n      i32 = ir.IntegerType.get_signless(32)","sourceCodeStart":1210,"sourceCodeEnd":1246,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L1210-L1246","documentation":"arrive_expect_tx uses the Hopper+ PTX instruction mbarrier.arrive.expect_tx, which reserves transaction-byte counting for TMA-style async transfers. On architectures with compute capability < 9.0 it raises NotImplementedError.","triggerScenarios":"barrier.arrive_expect_tx(bytes=N) on Ampere/Turing or any GPU where get_arch().major < 9; TMA async-copy pipelines run on pre-Hopper hardware.","commonSituations":"Running a TMA-based Mosaic kernel written for H100 on A100; CI fleets with mixed GPU generations; relying on cp.async.bulk semantics that only exist on sm90+.","solutions":["Target only sm90+ (H100/H200/B100) for expect_tx pipelines; guard with get_arch().major >= 9","Fall back to cp.async-based transfers with plain arrive/complete on older GPUs","Set the platform/arch explicitly in the kernel launch config so get_arch() matches the target device"],"exampleFix":"# before\nbarrier.arrive_expect_tx(bytes=chunk_bytes)\n# after\nfrom jax.experimental.mosaic.gpu.utils import get_arch\nif get_arch().major >= 9:\n    barrier.arrive_expect_tx(bytes=chunk_bytes)\nelse:\n    barrier.arrive()  # legacy cp.async path","handlingStrategy":"fallback","validationCode":"from jax.experimental.mosaic.gpu.utils import get_arch\nhopper_plus = get_arch().major >= 9","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Guard all expect_tx/complete_tx paths with get_arch().major >= 9","Provide a cp.async fallback path for pre-Hopper GPUs"],"tags":["jax","mosaic-gpu","barrier","tma","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"}