{"record":{"id":"b1d977739e8f5e0c","repo":"jax-ml/jax","slug":"predicate-not-supported-for-no-complete-arrive","errorCode":null,"errorMessage":"Predicate not supported for no-complete arrive","messagePattern":"Predicate not supported for no-complete arrive","errorType":"validation","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":1218,"sourceCode":"        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\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    )","sourceCodeStart":1200,"sourceCodeEnd":1236,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L1200-L1236","documentation":"arrive with can_complete=False lowers to nvvm.mbarrier_arrive_nocomplete, which has no predicate operand in the NVVM API. Passing a predicate together with the no-complete variant raises NotImplementedError.","triggerScenarios":"barrier.arrive(..., can_complete=False, predicate=p) — any combination of a non-None predicate and no-complete arrive.","commonSituations":"Adapting predicated arrival (used for single-lane arrival on Hopper) while also disabling completion counting; refactoring arrive calls so can_complete defaults changed under an existing predicate argument.","solutions":["Drop the predicate when can_complete=False; guard arrival upstream with when(predicate) around the call","Use the completing arrive (default) which supports predicates via inline asm","Restructure so only the designated lane calls arrive and skip predication entirely"],"exampleFix":"# before\nbarrier.arrive(can_complete=False, predicate=is_leader)\n# after\nwith when(is_leader):\n    barrier.arrive(can_complete=False)","handlingStrategy":"validation","validationCode":"assert not (predicate is not None and not can_complete), 'predicate requires completing arrive'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Wrap no-complete arrives in when(predicate) instead of passing a predicate arg","Keep can_complete/predicate combinations covered by unit tests"],"tags":["jax","mosaic-gpu","barrier","predicate","not-implemented"],"backgroundTag":"unsupported-argument-combination","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}