sgl-project/sglang · error · ValueError

kv-canary: {name} length must be >= {minimum}, got {actual}

Error message

kv-canary: {name} length must be >= {minimum}, got {actual}

What it means

Some inputs (like the LUT) only need a minimum length rather than an exact one; this error fires when such a tensor is shorter than the required minimum, meaning the kernel could read past its end.

Source

Thrown at python/sglang/kernels/ops/kv_canary/plan/utils.py:57

def _require_2d(tensor: torch.Tensor, name: str) -> None:
    if tensor.ndim != 2:
        raise ValueError(
            f"kv-canary: {name} must be 2-D, got shape {tuple(tensor.shape)}"
        )


def _require_len(tensor: torch.Tensor, name: str, expected: int) -> None:
    _require_1d(tensor=tensor, name=name)
    actual = int(tensor.shape[0])
    if actual != expected:
        raise ValueError(f"kv-canary: {name} length must be {expected}, got {actual}")


def _require_min_len(tensor: torch.Tensor, name: str, minimum: int) -> None:
    _require_1d(tensor=tensor, name=name)
    actual = int(tensor.shape[0])
    if actual < minimum:
        raise ValueError(f"kv-canary: {name} length must be >= {minimum}, got {actual}")


def _require_same_device(
    reference: torch.Tensor,
    reference_name: str,
    tensors: tuple[tuple[torch.Tensor, str], ...],
) -> None:
    for tensor, name in tensors:
        if tensor.device != reference.device:
            raise ValueError(
                f"kv-canary: {name} must be on {reference_name}'s device "
                f"{reference.device}, got {tensor.device}"
            )


@triton.jit
def _compute_window_start(prefix_lens, SWA_WINDOW: tl.constexpr):
    """Per-req window start: max(prefix_lens - SWA_WINDOW, 0) when SWA, else 0.

View on GitHub (pinned to 0132848349)

Solutions

1. Grow the LUT tensor to at least lut_len elements
2. Or correct lut_len to the tensor's real length when the config value is stale

Example fix

// before
lut_len = max_token_id + 1
lut_tensor = torch.zeros(max_token_id, dtype=torch.int64, device=dev)
// after
lut_len = max_token_id + 1
lut_tensor = torch.zeros(lut_len, dtype=torch.int64, device=dev)

assert lut_tensor.shape[0] >= max(lut_len, 1)

When it happens

Trigger: Calling launch_plan_offsets_kernel where lut_tensor (or another min-length-checked input) has shape[0] below the declared lut_len minimum — e.g. lut_len set larger than the actual LUT tensor.

Common situations: lut_len taken from config (max index + 1) while the allocated LUT is smaller; LUT truncated after a pool resize.

Related errors

• kv-canary: lut_len must be non-negative, got {lut_len}
• kv-canary: lut_len must be positive when has_swa_lut is True
• kv-canary: lut_len must be 0 when has_swa_lut is False
• kv-canary: {name} length must be {expected}, got {actual}
• g and beta must cover every q token

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