{"record":{"id":"03d183e694537ca1","repo":"sgl-project/sglang","slug":"kv-dtype-mismatch-kv-kv-dtype-q-q-dtype","errorCode":null,"errorMessage":"kv dtype mismatch: kv={kv.dtype}, q={q.dtype}","messagePattern":"kv dtype mismatch: kv=(.+?), q=(.+?)","errorType":"validation","errorClass":"RuntimeError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_prefill.py","lineNumber":241,"sourceCode":"    kv_indices_extend: torch.Tensor,\n    kv_indptr_extend: torch.Tensor,\n    attn_sink: torch.Tensor,\n    softmax_scale: float,\n) -> torch.Tensor:\n    if not q.is_cuda:\n        raise RuntimeError(\n            \"Triton sparse_attn_v4_paged_prefill requires CUDA/HIP tensors\"\n        )\n    if q.dtype not in (torch.bfloat16, torch.float16):\n        raise RuntimeError(\n            f\"sparse_attn_v4_paged_prefill expects fp16/bf16 q, got {q.dtype}\"\n        )\n    if unified_kv.dtype != q.dtype:\n        raise RuntimeError(\n            f\"unified_kv dtype mismatch: kv={unified_kv.dtype}, q={q.dtype}\"\n        )\n    if kv.dtype != q.dtype:\n        raise RuntimeError(f\"kv dtype mismatch: kv={kv.dtype}, q={q.dtype}\")\n    if unified_kv.size(-1) != kv.size(-1):\n        raise RuntimeError(\n            f\"head_dim mismatch: unified_kv={unified_kv.size(-1)}, kv={kv.size(-1)}\"\n        )\n\n    T, H, D = q.shape\n    out = torch.empty_like(q)\n    kv_indices_prefix = kv_indices_prefix.to(torch.int32).contiguous()\n    kv_indptr_prefix = kv_indptr_prefix.to(torch.int32).contiguous()\n    kv_indices_extend = kv_indices_extend.to(torch.int32).contiguous()\n    kv_indptr_extend = kv_indptr_extend.to(torch.int32).contiguous()\n\n    block_h = 16  # AMD MFMA min tile\n    block_d = triton.next_power_of_2(D)\n    block_k = 16 if D >= 256 else 32\n    _sparse_attn_v4_paged_prefill_kernel[(T, triton.cdiv(H, block_h))](\n        q,\n        unified_kv,","sourceCodeStart":223,"sourceCodeEnd":259,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_prefill.py#L223-L259","documentation":"sparse_attn_v4_paged_prefill takes both a paged unified_kv cache and a non-paged extend kv tensor (the current chunk's keys/values). The extend kv must match q's dtype exactly, same as the paged cache.","triggerScenarios":"Calling sparse_attn_v4_paged_prefill where the extend kv tensor (current chunk K/V) has a different dtype from q — e.g. extend path producing fp32 projections while q was cast to bf16.","commonSituations":"Extend/prefill projection layer left in fp32 while the rest is bf16; per-layer dtype overrides; a partial quantization setup where only some projections are cast.","solutions":["Cast kv to q.dtype before the call","Ensure the K/V projection outputs use the model's compute dtype consistently","Audit where kv is produced and cast at the source rather than at the kernel boundary"],"exampleFix":"// before\nout = sparse_attn_v4_paged_prefill(q, kv_fp32, ...)\n// after\nout = sparse_attn_v4_paged_prefill(q, kv_fp32.to(q.dtype), ...)","handlingStrategy":"validation","validationCode":"kv = kv.to(q.dtype)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Cast K/V projections to the compute dtype where they are produced","Avoid per-layer dtype overrides on attention inputs"],"tags":["attention","dtype","extend","kv"],"backgroundTag":"tensor-dtype-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}