sgl-project/sglang · error · ValueError

kv-canary: scatter_req_token_ids bs+1={bs + 1} exceeds BATCH

Error message

kv-canary: scatter_req_token_ids bs+1={bs + 1} exceeds BATCH_BLOCK={_SCATTER_BATCH_BLOCK}; bump _SCATTER_BATCH_BLOCK if real workloads need this

What it means

The Triton scatter kernel keeps the whole offsets vector in registers via a single BLOCK of size _SCATTER_BATCH_BLOCK; if bs+1 exceeds that compile-time constant the kernel would silently drop requests, so the launcher raises ValueError instead. The message says to bump the constant if real workloads need larger batches.

Source

Thrown at python/sglang/kernels/ops/kv_canary/scatter_req_token_ids.py:93

    if req_pool_indices.dtype != torch.int64:
        raise TypeError(
            f"kv-canary: scatter_req_token_ids req_pool_indices must be int64, got "
            f"{req_pool_indices.dtype}"
        )
    if pool_out.dtype != torch.int32:
        raise TypeError(
            f"kv-canary: scatter_req_token_ids pool_out must be int32, got "
            f"{pool_out.dtype}"
        )

    bs = int(req_pool_indices.shape[0])
    if int(offsets.shape[0]) != bs + 1:
        raise ValueError(
            f"kv-canary: scatter_req_token_ids offsets length {offsets.shape[0]} != "
            f"bs+1 ({bs + 1})"
        )
    if bs + 1 > _SCATTER_BATCH_BLOCK:
        raise ValueError(
            f"kv-canary: scatter_req_token_ids bs+1={bs + 1} exceeds BATCH_BLOCK="
            f"{_SCATTER_BATCH_BLOCK}; bump _SCATTER_BATCH_BLOCK if real workloads need this"
        )

    num_tokens = int(flat_in.shape[0])
    if num_tokens == 0:
        return

    pool_stride0 = int(pool_out.stride(0))
    pool_max_context_len = int(pool_out.shape[1])

    grid = (triton.cdiv(num_tokens, _SCATTER_TOKEN_BLOCK),)
    _scatter_req_token_ids_kernel[grid](
        flat_in,
        offsets,
        req_pool_indices,
        pool_out,
        num_tokens=num_tokens,

View on GitHub (pinned to 0132848349)

Solutions

1. Split the batch into chunks of at most _SCATTER_BATCH_BLOCK - 1 requests and launch per chunk with sliced offsets/indices
2. Bump _SCATTER_BATCH_BLOCK in scatter_req_token_ids.py and recompile if your workload genuinely needs bigger single-launch batches
3. Check the current constant first to know your effective max bs

Example fix

# before
launch_scatter(..., offsets=offsets, req_pool_indices=rp)  # bs too big
# after
CHUNK = _SCATTER_BATCH_BLOCK - 1
for i in range(0, bs, CHUNK):
    launch_scatter(..., offsets=offsets[i:i+CHUNK+1], req_pool_indices=rp[i:i+CHUNK])

bs = req_pool_indices.shape[0]
assert bs + 1 <= _SCATTER_BATCH_BLOCK, f"bs={bs} exceeds kernel limit"

When it happens

Trigger: Calling launch_scatter_req_token_ids_kernel with a batch size larger than _SCATTER_BATCH_BLOCK - 1 (the kernel was sized for 'bs <= a few thousand').

Common situations: Scaling up server concurrency so the running batch exceeds the kernel's design limit; running large batch microbenchmarks.

Related errors

• kv_scales supplied but unified_kv is {unified_kv.dtype}, exp
• unified_kv dtype mismatch: kv={unified_kv.dtype}, q={q.dtype
• kv-canary: {name} must be contiguous
• v_cache must be provided
• k_cache can only be None when only_qv=True

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/dbca6d79092499b7. Report an issue: GitHub.