{"record":{"id":"ab8d20bf7261083d","repo":"jax-ml/jax","slug":"unsupported-scope-scope-ab8d20","errorCode":null,"errorMessage":"Unsupported scope: {scope}","messagePattern":"Unsupported scope: (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/experimental/mosaic/gpu/utils.py","lineNumber":1120,"sourceCode":"    ) -> ir.Value:\n    i1 = ir.IntegerType.get_signless(1)\n    i32 = ir.IntegerType.get_signless(32)\n    parity = arith.extui(i32, parity)\n    wait_complete = nvvm.mbarrier_test_wait(self.get_ptr(), parity)\n\n    if scope == ThreadSubset.WARPGROUP:\n      wait_complete = llvm.inline_asm(\n          i1,\n          [warpgroup_barrier_idx(sync=False), wait_complete],\n          f\"bar.red.or.pred $0, $1, {WARPGROUP_SIZE}, $2;\",\n          \"=b,r,b\",\n          has_side_effects=True,\n      )\n      wait_complete = cast(ir.OpResult[ir.IntegerType], wait_complete)\n    elif scope == ThreadSubset.WARP:\n      wait_complete = nvvm.vote_sync(c(0xFFFFFFFF, i32), wait_complete, \"any\")\n    else:\n      raise ValueError(f\"Unsupported scope: {scope}\")\n\n    if orders_tensor_core:\n      with when(wait_complete):\n        nvvm.tcgen05_fence(nvvm.Tcgen05FenceKind.AFTER_THREAD_SYNC)\n    return wait_complete\n\n  def test(\n      self,\n      orders_tensor_core: bool = False,\n      scope: ThreadSubset = ThreadSubset.WARPGROUP,\n  ) -> ir.Value:\n    parities = memref.load(self.phases, [])\n    parity, new_parities = self.update_parities(parities)\n    wait_complete = self.test_parity(parity, orders_tensor_core, scope)\n    with when(wait_complete):\n      memref.store(new_parities, self.phases, [])\n    return wait_complete\n","sourceCodeStart":1102,"sourceCodeEnd":1138,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/experimental/mosaic/gpu/utils.py#L1102-L1138","documentation":"The barrier test/parity helper only knows how to broadcast completion for BLOCK and WARP thread scopes. Any other ThreadSubset (e.g. WARPGROUP or an invalid value) reaches the else branch and raises ValueError('Unsupported scope').","triggerScenarios":"Calling the barrier test helper with scope=ThreadSubset.WARPGROUP or a raw/None scope value; enum value changed by an upstream refactor so it no longer matches the handled cases.","commonSituations":"Adding new scopes to ThreadSubset without updating this dispatch; passing a default parameter left as an unhandled enum member; version drift after upgrading jax where scope handling moved.","solutions":["Pass ThreadSubset.BLOCK or ThreadSubset.WARP explicitly","For warpgroup-wide completion, wrap a WARP-scope test plus warpgroup_barrier()","Check the installed jax version's ThreadSubset members and update call sites"],"exampleFix":"# before\nok = barrier.test(parity, scope=ThreadSubset.WARPGROUP)\n# after\nok = barrier.test(parity, scope=ThreadSubset.WARP)\nutils.warpgroup_barrier()  # extend to warpgroup","handlingStrategy":"validation","validationCode":"assert scope in (ThreadSubset.BLOCK, ThreadSubset.WARP), scope","typeGuard":"def supported_test_scope(s) -> bool:\n    return s in (ThreadSubset.BLOCK, ThreadSubset.WARP)","tryCatchPattern":null,"preventionTips":["Check the ThreadSubset members supported by each barrier API before parameterizing scope","Wrap warpgroup needs as WARP test + warpgroup_barrier()"],"tags":["jax","mosaic-gpu","barrier","scope","invalid-argument"],"backgroundTag":"unsupported-enum-value","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}