{"record":{"id":"40eac09a593ee127","repo":"sgl-project/sglang","slug":"q-k-and-v-must-have-the-same-device-and-dtype","errorCode":null,"errorMessage":"q, k, and v must have the same device and dtype","messagePattern":"q, k, and v must have the same device and dtype","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py","lineNumber":68,"sourceCode":"    tl.store(output_ptr + output_base, q, mask=mask)\n    tl.store(output_ptr + output_base + head_size, k, mask=mask)\n    tl.store(output_ptr + output_base + 2 * head_size, v, mask=mask)\n\n\ndef pack_qkv_destination_major(\n    q: torch.Tensor,\n    k: torch.Tensor,\n    v: torch.Tensor,\n    world_size: int,\n    out: torch.Tensor | None = None,\n) -> torch.Tensor:\n    \"\"\"Pack matching ``[rows, global_heads, head_size]`` Q/K/V tensors.\"\"\"\n    if q.dim() != 3 or q.shape != k.shape or q.shape != v.shape:\n        raise ValueError(\"q, k, and v must have the same 3D shape\")\n    if not (q.is_cuda and k.is_cuda and v.is_cuda):\n        raise ValueError(\"q, k, and v must be CUDA tensors\")\n    if not (q.device == k.device == v.device and q.dtype == k.dtype == v.dtype):\n        raise ValueError(\"q, k, and v must have the same device and dtype\")\n    if q.stride(-1) != 1 or k.stride(-1) != 1 or v.stride(-1) != 1:\n        raise ValueError(\"q, k, and v must be contiguous in head_size\")\n    if world_size < 1 or q.shape[1] % world_size != 0:\n        raise ValueError(\"world_size must be positive and divide global_heads\")\n\n    rows, global_heads, head_size = q.shape\n    local_heads = global_heads // world_size\n    expected_shape = (world_size, rows, local_heads, 3 * head_size)\n    if out is not None:\n        if not (\n            out.shape == expected_shape\n            and out.is_contiguous()\n            and out.dtype == q.dtype\n            and out.device == q.device\n        ):\n            raise ValueError(\n                \"out must be a contiguous tensor with the expected shape, \"\n                \"device, and dtype\"","sourceCodeStart":50,"sourceCodeEnd":86,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/layout/ulysses_qkv_triton.py#L50-L86","documentation":"Before packing, the kernel requires q, k, v to live on the same CUDA device and share one dtype. Mixed devices (e.g. cuda:0 vs cuda:1) or mixed dtypes (fp16 q with bf16 k) raise this ValueError.","triggerScenarios":"Passing tensors allocated on different GPUs in multi-GPU setups without proper device placement, or tensors cast to different precisions (e.g. q in bf16 but v still in fp16) during half-precision conversion of a model.","commonSituations":"Tensor-parallel or multi-GPU diffusion serving where per-rank tensors land on different devices; partially-applied dtype conversions (model.to(torch.bfloat16) missing some buffers); legacy fp16 checkpoints mixed with bf16 activations.","solutions":["Unify dtype: cast k and v (or all three) to a single dtype with .to(q.dtype)","Verify q.device == k.device == v.device; move tensors with .to(q.device)","Check that model weights and activations were converted consistently (model.to(dtype))","In TP setups, ensure the rank's tensors are on its assigned device before packing"],"exampleFix":"# before\npacked = pack_qkv_destination_major(q_bf16, k_fp16, v_bf16, ws)\n# after\nk = k.to(q.dtype)\npacked = pack_qkv_destination_major(q, k, v, ws)","handlingStrategy":"validation","validationCode":"assert q.device == k.device == v.device and q.dtype == k.dtype == v.dtype\nk = k.to(q.dtype); v = v.to(q.dtype)","typeGuard":"def same_device_dtype(q, k, v) -> bool:\n    return q.device == k.device == v.device and q.dtype == k.dtype == v.dtype","tryCatchPattern":null,"preventionTips":["Convert whole model with model.to(dtype) once","Pin per-rank device in TP setups"],"tags":["dtype","device","mismatch","ulysses"],"backgroundTag":"dtype-device-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}