sgl-project/sglang · error · ValueError

QKV tensors must be CUDA bfloat16 tensors

Error message

QKV tensors must be CUDA bfloat16 tensors

What it means

The Hunyuan QKV pack Triton kernel requires QKV tensors on CUDA with dtype bfloat16. Cast/move them accordingly.

Solutions

1. Cast all six tensors to bfloat16 and .cuda() before the call
2. Run the model with bfloat16 dtype (torch_dtype=torch.bfloat16)
3. Use a non-fused RoPE path for other precisions

Example fix

# before
out = hunyuan_qkv_rope_pack(img_q, ...)  # fp16 tensors
# after
out = hunyuan_qkv_rope_pack(img_q.to(torch.bfloat16).cuda(), img_k.to(torch.bfloat16).cuda(), ...)

if not all(t.is_cuda and t.dtype == torch.bfloat16 for t in tensors):
    tensors = [t.to(torch.bfloat16).cuda() for t in tensors]

def is_cuda_bf16(t: torch.Tensor) -> bool:
    return t.is_cuda and t.dtype == torch.bfloat16

When it happens

Trigger: Passing fp16 or fp32 QKV tensors, or CPU tensors not yet moved to GPU, to hunyuan_qkv_rope_pack.

Common situations: Running the model in float16 precision mode, or under a CPU meta/init pass where projections are still on CPU; mixed-precision configs that leave attention in fp32.

Related errors

• QKV tensors must be on the same CUDA device
• Unsupported interleaved_rope_fp64 dtype
• Unsupported ltx25_decoder_rope dtype
• f"Sol-Attn requires bfloat16 activations, got
• Fused QK-Norm + RoPE kernel only supports float16/bfloat16…

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

Appendix: source

Thrown at python/sglang/kernels/ops/diffusion/rope/hunyuan_qkv_pack_triton.py:165

    tl.store(output_ptr + 2 * plane_stride + output_row + even[None, :], v0, mask=mask)
    tl.store(output_ptr + 2 * plane_stride + output_row + odd[None, :], v1, mask=mask)


def hunyuan_qkv_rope_pack(
    img_q: torch.Tensor,
    img_k: torch.Tensor,
    img_v: torch.Tensor,
    txt_q: torch.Tensor,
    txt_k: torch.Tensor,
    txt_v: torch.Tensor,
    cos: torch.Tensor,
    sin: torch.Tensor,
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
    tensors = (img_q, img_k, img_v, txt_q, txt_k, txt_v)
    if any(x.ndim != 4 for x in tensors):
        raise ValueError("QKV tensors must have shape [B, S, H, D]")
    if any(not x.is_cuda or x.dtype != torch.bfloat16 for x in tensors):
        raise ValueError("QKV tensors must be CUDA bfloat16 tensors")
    if any(x.device != img_q.device for x in tensors):
        raise ValueError("QKV tensors must be on the same CUDA device")
    batch, img_tokens, num_heads, head_dim = img_q.shape
    txt_tokens = txt_q.shape[1]
    expected_img = (batch, img_tokens, num_heads, head_dim)
    expected_txt = (batch, txt_tokens, num_heads, head_dim)
    if any(tuple(x.shape) != expected_img for x in (img_q, img_k, img_v)):
        raise ValueError("image QKV shapes must match")
    if any(tuple(x.shape) != expected_txt for x in (txt_q, txt_k, txt_v)):
        raise ValueError("text QKV shapes must match")
    if any(x.stride(-1) != 1 for x in tensors):
        raise ValueError("QKV last dimensions must be contiguous")
    if head_dim <= 0 or head_dim > 128 or head_dim % 2:
        raise ValueError("head_dim must be positive, even, and <= 128")
    if cos.ndim != 2 or sin.ndim != 2 or cos.shape != sin.shape:
        raise ValueError("cos and sin must have matching [S, D/2] shapes")
    if cos.shape[0] < img_tokens or cos.shape[1] != head_dim // 2:
        raise ValueError("cos/sin shape does not cover image tokens and head_dim")

View on GitHub (pinned to 0132848349)