sgl-project/sglang · error · RuntimeError

kv dtype mismatch: kv={kv.dtype}, q={q.dtype}

Error message

kv dtype mismatch: kv={kv.dtype}, q={q.dtype}

What it means

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.

Source

Thrown at python/sglang/kernels/ops/attention/dsv4/unified_kv_kernels/paged_prefill.py:241

    kv_indices_extend: torch.Tensor,
    kv_indptr_extend: torch.Tensor,
    attn_sink: torch.Tensor,
    softmax_scale: float,
) -> torch.Tensor:
    if not q.is_cuda:
        raise RuntimeError(
            "Triton sparse_attn_v4_paged_prefill requires CUDA/HIP tensors"
        )
    if q.dtype not in (torch.bfloat16, torch.float16):
        raise RuntimeError(
            f"sparse_attn_v4_paged_prefill expects fp16/bf16 q, got {q.dtype}"
        )
    if unified_kv.dtype != q.dtype:
        raise RuntimeError(
            f"unified_kv dtype mismatch: kv={unified_kv.dtype}, q={q.dtype}"
        )
    if kv.dtype != q.dtype:
        raise RuntimeError(f"kv dtype mismatch: kv={kv.dtype}, q={q.dtype}")
    if unified_kv.size(-1) != kv.size(-1):
        raise RuntimeError(
            f"head_dim mismatch: unified_kv={unified_kv.size(-1)}, kv={kv.size(-1)}"
        )

    T, H, D = q.shape
    out = torch.empty_like(q)
    kv_indices_prefix = kv_indices_prefix.to(torch.int32).contiguous()
    kv_indptr_prefix = kv_indptr_prefix.to(torch.int32).contiguous()
    kv_indices_extend = kv_indices_extend.to(torch.int32).contiguous()
    kv_indptr_extend = kv_indptr_extend.to(torch.int32).contiguous()

    block_h = 16  # AMD MFMA min tile
    block_d = triton.next_power_of_2(D)
    block_k = 16 if D >= 256 else 32
    _sparse_attn_v4_paged_prefill_kernel[(T, triton.cdiv(H, block_h))](
        q,
        unified_kv,

View on GitHub (pinned to 0132848349)

Solutions

1. Cast kv to q.dtype before the call
2. Ensure the K/V projection outputs use the model's compute dtype consistently
3. Audit where kv is produced and cast at the source rather than at the kernel boundary

Example fix

// before
out = sparse_attn_v4_paged_prefill(q, kv_fp32, ...)
// after
out = sparse_attn_v4_paged_prefill(q, kv_fp32.to(q.dtype), ...)

kv = kv.to(q.dtype)

When it happens

Trigger: 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.

Common situations: 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.

Related errors

• sparse_attn_v4_paged_decode expects fp16/bf16 q, got {q.dtyp
• unified_kv dtype mismatch: kv={unified_kv.dtype}, q={q.dtype
• sparse_attn_v4_paged_prefill expects fp16/bf16 q, got {q.dty
• unified_kv dtype mismatch: kv={unified_kv.dtype}, q={q.dtype
• {name}_block tensors must have dtype torch.int32

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