sgl-project/sglang · error · ValueError

kv-canary: {name} must be on {reference_name}'s device {refe

Error message

kv-canary: {name} must be on {reference_name}'s device {reference.device}, got {tensor.device}

What it means

A helper in the kv-canary plan utilities enforces that every tensor passed alongside a reference tensor lives on the same torch device as the reference. If any tensor's device differs, a ValueError is raised naming the offending tensor and the expected device. This guards the Triton offsets kernel, which assumes all inputs are co-located (typically all on CPU for the host-side plan path).

Source

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

    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.
    Works for tile and scalar inputs (broadcasts via prefix_lens shape).
    """
    if SWA_WINDOW > 0:
        clipped = prefix_lens - SWA_WINDOW
        return tl.where(clipped > 0, clipped, 0)
    else:
        return prefix_lens - prefix_lens


@triton.jit

View on GitHub (pinned to 0132848349)

Solutions

1. Move all tensors to the same device as the reference tensor before calling the API (e.g. t.to(reference.device))
2. Check tensor.device for each input in the caller and normalize them
3. Ensure your scheduler does not slice inputs from pools on a different device than the plan tensors

Example fix

# before
launch(..., offsets=offsets_cpu, req_pool_indices=req_indices_cuda)
# after
req_indices = req_indices_cuda.to(offsets_cpu.device)
launch(..., offsets=offsets_cpu, req_pool_indices=req_indices)

devices = {t.device for t in (ref, *others)}
assert len(devices) == 1, f"device mix: {devices}"

def all_same_device(ref: torch.Tensor, *ts: torch.Tensor) -> bool:
    return all(t.device == ref.device for t in ts)

When it happens

Trigger: Calling _validate_offsets_kernel_inputs (via the plan API) with e.g. offsets on CPU but req_pool_indices on cuda:0, or any tensor pair where one was moved with .to('cuda') and another left on CPU.

Common situations: Mixing host-side planning tensors with GPU-resident scheduler tensors; passing a tensor produced by torch.zeros (CPU default) next to one sliced off a CUDA allocation pool.

Related errors

• v_cache must be provided
• kv-canary: launch_canary_plan_kernels_torch_reference verify
• kv-canary: launch_canary_plan_kernels_torch_reference requir
• kv-canary: scatter_req_token_ids flat_in must be int64, got
• kv-canary: scatter_req_token_ids offsets must be int64, got

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