{"record":{"id":"352cf38b3c430a02","repo":"sgl-project/sglang","slug":"qkv-tensors-must-have-shape-b-s-h-d","errorCode":null,"errorMessage":"QKV tensors must have shape [B, S, H, D]","messagePattern":"QKV tensors must have shape \\[B, S, H, D\\]","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/diffusion/rope/hunyuan_qkv_pack_triton.py","lineNumber":163,"sourceCode":"    tl.store(output_ptr + plane_stride + output_row + even[None, :], ok0, mask=mask)\n    tl.store(output_ptr + plane_stride + output_row + odd[None, :], ok1, mask=mask)\n    tl.store(output_ptr + 2 * plane_stride + output_row + even[None, :], v0, mask=mask)\n    tl.store(output_ptr + 2 * plane_stride + output_row + odd[None, :], v1, mask=mask)\n\n\ndef hunyuan_qkv_rope_pack(\n    img_q: torch.Tensor,\n    img_k: torch.Tensor,\n    img_v: torch.Tensor,\n    txt_q: torch.Tensor,\n    txt_k: torch.Tensor,\n    txt_v: torch.Tensor,\n    cos: torch.Tensor,\n    sin: torch.Tensor,\n) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:\n    tensors = (img_q, img_k, img_v, txt_q, txt_k, txt_v)\n    if any(x.ndim != 4 for x in tensors):\n        raise ValueError(\"QKV tensors must have shape [B, S, H, D]\")\n    if any(not x.is_cuda or x.dtype != torch.bfloat16 for x in tensors):\n        raise ValueError(\"QKV tensors must be CUDA bfloat16 tensors\")\n    if any(x.device != img_q.device for x in tensors):\n        raise ValueError(\"QKV tensors must be on the same CUDA device\")\n    batch, img_tokens, num_heads, head_dim = img_q.shape\n    txt_tokens = txt_q.shape[1]\n    expected_img = (batch, img_tokens, num_heads, head_dim)\n    expected_txt = (batch, txt_tokens, num_heads, head_dim)\n    if any(tuple(x.shape) != expected_img for x in (img_q, img_k, img_v)):\n        raise ValueError(\"image QKV shapes must match\")\n    if any(tuple(x.shape) != expected_txt for x in (txt_q, txt_k, txt_v)):\n        raise ValueError(\"text QKV shapes must match\")\n    if any(x.stride(-1) != 1 for x in tensors):\n        raise ValueError(\"QKV last dimensions must be contiguous\")\n    if head_dim <= 0 or head_dim > 128 or head_dim % 2:\n        raise ValueError(\"head_dim must be positive, even, and <= 128\")\n    if cos.ndim != 2 or sin.ndim != 2 or cos.shape != sin.shape:\n        raise ValueError(\"cos and sin must have matching [S, D/2] shapes\")","sourceCodeStart":145,"sourceCodeEnd":181,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/diffusion/rope/hunyuan_qkv_pack_triton.py#L145-L181","documentation":"hunyuan_qkv_rope_pack fuses RoPE with QKV packing for Hunyuan video diffusion and requires all six img/txt Q/K/V tensors to be 4D (B, S, H, D). Anything else can't be packed by the Triton kernel.","triggerScenarios":"Calling hunyuan_qkv_rope_pack with any of img_q/img_k/img_v/txt_q/txt_k/txt_v having ndim != 4 (e.g. 3D (B, S, H*D) unreshaped projections).","commonSituations":"Forgetting to reshape attention projections from (B, S, H*D) to (B, S, H, D) before the fused RoPE/pack call.","solutions":["Reshape each projection: q.view(B, S, num_heads, head_dim) before calling","Check upstream linear layer outputs are reshaped consistently for all six tensors"],"exampleFix":"# before\nimg_q, img_k, img_v = proj(x).chunk(3, dim=-1)  # (B,S,H*D)\nout = hunyuan_qkv_rope_pack(img_q, ...)\n# after\nimg_q = img_q.view(B, S, H, D)\nimg_k = img_k.view(B, S, H, D)\nimg_v = img_v.view(B, S, H, D)\nout = hunyuan_qkv_rope_pack(img_q, ...)","handlingStrategy":"validation","validationCode":"assert all(t.ndim == 4 for t in (img_q, img_k, img_v, txt_q, txt_k, txt_v))","typeGuard":"def is_4d_bshd(t: torch.Tensor) -> bool:\n    return t.ndim == 4","tryCatchPattern":null,"preventionTips":["Reshape projections to (B, S, H, D) immediately after the QKV linear"],"tags":["shape","rope","attention","hunyuan","diffusion"],"backgroundTag":"tensor-shape-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}