{"record":{"id":"bc66adb90d58fae9","repo":"sgl-project/sglang","slug":"kv-canary-name-must-be-on-reference-name-s-de","errorCode":null,"errorMessage":"kv-canary: {name} must be on {reference_name}'s device {reference.device}, got {tensor.device}","messagePattern":"kv-canary: (.+?) must be on (.+?)'s device (.+?), got (.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/kv_canary/plan/utils.py","lineNumber":67,"sourceCode":"    if actual != expected:\n        raise ValueError(f\"kv-canary: {name} length must be {expected}, got {actual}\")\n\n\ndef _require_min_len(tensor: torch.Tensor, name: str, minimum: int) -> None:\n    _require_1d(tensor=tensor, name=name)\n    actual = int(tensor.shape[0])\n    if actual < minimum:\n        raise ValueError(f\"kv-canary: {name} length must be >= {minimum}, got {actual}\")\n\n\ndef _require_same_device(\n    reference: torch.Tensor,\n    reference_name: str,\n    tensors: tuple[tuple[torch.Tensor, str], ...],\n) -> None:\n    for tensor, name in tensors:\n        if tensor.device != reference.device:\n            raise ValueError(\n                f\"kv-canary: {name} must be on {reference_name}'s device \"\n                f\"{reference.device}, got {tensor.device}\"\n            )\n\n\n@triton.jit\ndef _compute_window_start(prefix_lens, SWA_WINDOW: tl.constexpr):\n    \"\"\"Per-req window start: max(prefix_lens - SWA_WINDOW, 0) when SWA, else 0.\n    Works for tile and scalar inputs (broadcasts via prefix_lens shape).\n    \"\"\"\n    if SWA_WINDOW > 0:\n        clipped = prefix_lens - SWA_WINDOW\n        return tl.where(clipped > 0, clipped, 0)\n    else:\n        return prefix_lens - prefix_lens\n\n\n@triton.jit","sourceCodeStart":49,"sourceCodeEnd":85,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/kv_canary/plan/utils.py#L49-L85","documentation":"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).","triggerScenarios":"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.","commonSituations":"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.","solutions":["Move all tensors to the same device as the reference tensor before calling the API (e.g. t.to(reference.device))","Check tensor.device for each input in the caller and normalize them","Ensure your scheduler does not slice inputs from pools on a different device than the plan tensors"],"exampleFix":"# before\nlaunch(..., offsets=offsets_cpu, req_pool_indices=req_indices_cuda)\n# after\nreq_indices = req_indices_cuda.to(offsets_cpu.device)\nlaunch(..., offsets=offsets_cpu, req_pool_indices=req_indices)","handlingStrategy":"validation","validationCode":"devices = {t.device for t in (ref, *others)}\nassert len(devices) == 1, f\"device mix: {devices}\"","typeGuard":"def all_same_device(ref: torch.Tensor, *ts: torch.Tensor) -> bool:\n    return all(t.device == ref.device for t in ts)","tryCatchPattern":null,"preventionTips":["Normalize all plan/offset tensors to one device (usually CPU) in one place before calling kv-canary APIs","Assert device equality in tests since CPU/GPU mix-ups pass silently elsewhere"],"tags":["kv-cache","device-mismatch","torch","validation"],"backgroundTag":"tensor-device-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}