sgl-project/sglang · error · ValueError

kv-canary: scatter_req_token_ids offsets length {offsets.sha

Error message

kv-canary: scatter_req_token_ids offsets length {offsets.shape[0]} != bs+1 ({bs + 1})

What it means

offsets must be a CSR prefix vector of exactly length bs+1 (where bs = req_pool_indices.shape[0]): offsets[i]:offsets[i+1] is request i's span in flat_in. Any other length raises ValueError because the kernel reads bs+1 offsets per block.

Source

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

    if offsets.dtype != torch.int64:
        raise TypeError(
            f"kv-canary: scatter_req_token_ids offsets must be int64, got "
            f"{offsets.dtype}"
        )
    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](

View on GitHub (pinned to 0132848349)

Solutions

1. Build offsets as a length-bs+1 prefix sum: offsets = torch.zeros(bs+1, dtype=torch.int64); offsets[1:] = lens.cumsum(0)
2. Rebuild offsets whenever the batch composition changes; do not cache across scheduler iterations
3. Assert offsets.shape[0] == req_pool_indices.shape[0] + 1 before calling

Example fix

# before
offsets = lens  # lengths, not prefix offsets
# after
offsets = torch.zeros(lens.shape[0] + 1, dtype=torch.int64)
offsets[1:] = torch.cumsum(lens, 0)

assert offsets.shape[0] == req_pool_indices.shape[0] + 1, (offsets.shape, req_pool_indices.shape)
assert (torch.diff(offsets) >= 0).all()

When it happens

Trigger: Passing bs offsets (per-request lengths instead of prefix sums), or bs+2, or an offsets vector built for a different batch size than req_pool_indices.

Common situations: Off-by-one confusion between lengths and prefix offsets; stale offsets tensor from a previous batch after the batch size changed.

Related errors

• rope_pool_fused expects pool tensors to be 3-D
• k_pool has incompatible shape {k_pool.shape}
• v_pool shape must match k_pool shape, got {v_pool.shape} vs
• D={D_check} must be divisible by GROUP_SIZE={_FP8_GROUP_SIZE
• kv_scales shape {tuple(kv_scales.shape)} does not match expe

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