jax-ml/jax · error · ValueError

Unsupported scope: {scope}

Error message

Unsupported scope: {scope}

What it means

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').

Source

Thrown at jax/experimental/mosaic/gpu/utils.py:1120

    ) -> ir.Value:
    i1 = ir.IntegerType.get_signless(1)
    i32 = ir.IntegerType.get_signless(32)
    parity = arith.extui(i32, parity)
    wait_complete = nvvm.mbarrier_test_wait(self.get_ptr(), parity)

    if scope == ThreadSubset.WARPGROUP:
      wait_complete = llvm.inline_asm(
          i1,
          [warpgroup_barrier_idx(sync=False), wait_complete],
          f"bar.red.or.pred $0, $1, {WARPGROUP_SIZE}, $2;",
          "=b,r,b",
          has_side_effects=True,
      )
      wait_complete = cast(ir.OpResult[ir.IntegerType], wait_complete)
    elif scope == ThreadSubset.WARP:
      wait_complete = nvvm.vote_sync(c(0xFFFFFFFF, i32), wait_complete, "any")
    else:
      raise ValueError(f"Unsupported scope: {scope}")

    if orders_tensor_core:
      with when(wait_complete):
        nvvm.tcgen05_fence(nvvm.Tcgen05FenceKind.AFTER_THREAD_SYNC)
    return wait_complete

  def test(
      self,
      orders_tensor_core: bool = False,
      scope: ThreadSubset = ThreadSubset.WARPGROUP,
  ) -> ir.Value:
    parities = memref.load(self.phases, [])
    parity, new_parities = self.update_parities(parities)
    wait_complete = self.test_parity(parity, orders_tensor_core, scope)
    with when(wait_complete):
      memref.store(new_parities, self.phases, [])
    return wait_complete

View on GitHub (pinned to 1e1c6a8fc0)

Solutions

1. Pass ThreadSubset.BLOCK or ThreadSubset.WARP explicitly
2. For warpgroup-wide completion, wrap a WARP-scope test plus warpgroup_barrier()
3. Check the installed jax version's ThreadSubset members and update call sites

Example fix

# before
ok = barrier.test(parity, scope=ThreadSubset.WARPGROUP)
# after
ok = barrier.test(parity, scope=ThreadSubset.WARP)
utils.warpgroup_barrier()  # extend to warpgroup

assert scope in (ThreadSubset.BLOCK, ThreadSubset.WARP), scope

def supported_test_scope(s) -> bool:
    return s in (ThreadSubset.BLOCK, ThreadSubset.WARP)

When it happens

Trigger: 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.

Common situations: 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.

Related errors

• Can only await on CTA-local barriers
• Indexing barrier with {transforms} not supported in GPU inte
• Expected an `NDIndexer`, but got {transforms[0]}
• Attempting to operate on barrier without indexing, but `num_
• Expected a single barrier, got a barrier reference with shap

AI-assisted analysis of jax-ml/jax@1e1c6a8fc0 (2026-08-27). Data as JSON: /api/errors/ab8d20bf7261083d. Report an issue: GitHub.