{"record":{"id":"3f079a5dfcf31faa","repo":"sgl-project/sglang","slug":"invalid-packed-mixed-qkv-last-dim-qkv-dim-for","errorCode":null,"errorMessage":"Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V={V}.","messagePattern":"Invalid packed `mixed_qkv` last dim=(.+?) for HV=(.+?), V=(.+?)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/fla/fused_recurrent.py","lineNumber":344,"sourceCode":"        raise ValueError(\"`initial_state` must be contiguous in the last dim.\")\n    HV, V, K = initial_state.shape[-3:]\n    if a.shape[1] != HV or b.shape[1] != HV:\n        raise ValueError(\n            f\"`a`/`b` must have shape [B, HV] with HV={HV} (got a.shape={tuple(a.shape)}, b.shape={tuple(b.shape)}).\"\n        )\n    if A_log.numel() != HV or dt_bias.numel() != HV:\n        raise ValueError(\n            f\"`A_log` and `dt_bias` must have {HV} elements (got A_log.numel()={A_log.numel()}, dt_bias.numel()={dt_bias.numel()}).\"\n        )\n    if out.shape != (B, 1, HV, V):\n        raise ValueError(\n            f\"`out` must have shape {(B, 1, HV, V)} (got out.shape={tuple(out.shape)}).\"\n        )\n\n    qkv_dim = mixed_qkv.shape[1]\n    qk_dim = qkv_dim - HV * V\n    if qk_dim <= 0 or qk_dim % 2 != 0:\n        raise ValueError(\n            f\"Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V={V}.\"\n        )\n    q_dim = qk_dim // 2\n    if q_dim % K != 0:\n        raise ValueError(f\"Invalid packed Q size {q_dim}: must be divisible by K={K}.\")\n    H = q_dim // K\n    if H <= 0 or HV % H != 0:\n        raise ValueError(\n            f\"Invalid head config inferred from mixed_qkv: H={H}, HV={HV}.\"\n        )\n\n    BK = triton.next_power_of_2(K)\n    if triton.cdiv(K, BK) != 1:\n        raise ValueError(\n            f\"Packed decode kernel only supports NK=1 (got K={K}, BK={BK}).\"\n        )\n    BV = min(triton.next_power_of_2(V), 32)\n    num_stages = 3","sourceCodeStart":326,"sourceCodeEnd":362,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/fla/fused_recurrent.py#L326-L362","documentation":"Thrown by fused_recurrent_gated_delta_rule_packed_decode when the packed mixed_qkv tensor's last dimension cannot be split into Q/K and V parts. The kernel computes qk_dim = mixed_qkv.shape[1] - HV*V and requires it to be positive and even (Q and K have equal head dim). If the layout or head counts don't match the expected packed q|k|v format, this validation fails.","triggerScenarios":"Calling fused_recurrent_gated_delta_rule_packed_decode with a mixed_qkv tensor whose shape[1] is not HV*V + 2*H*K, e.g. wrong V/HV values, an extra gate column packed in, or a transposed/non-packed tensor.","commonSituations":"Model config mismatch (V, HV, K from initial_state not matching the projection output width), changing head_dim or num_v_heads without regenerating the packed qkv buffer, or packing additional tensors (e.g. gates) into mixed_qkv.","solutions":["Check that mixed_qkv.shape[1] == HV*V + 2*H*K for your model config and that qk_dim = shape[1] - HV*V is even and > 0","Verify the V, HV, K values inferred from initial_state.shape[-3:] match the model's projection layer sizes","Ensure mixed_qkv was built by concatenating q, k, v along dim 1 with no extra columns"],"exampleFix":"# before\nmixed_qkv = torch.cat([q, k, v, gate], dim=1)  # extra gate column\n# after\nmixed_qkv = torch.cat([q, k, v], dim=1)  # qk_dim = 2*H*K, even","handlingStrategy":"validation","validationCode":"qkv_dim = mixed_qkv.shape[1]\nqk_dim = qkv_dim - HV * V\nassert qk_dim > 0 and qk_dim % 2 == 0, (qkv_dim, HV, V)","typeGuard":"def valid_packed_qkv(mixed_qkv, HV, V, K):\n    qk = mixed_qkv.shape[1] - HV * V\n    return mixed_qkv.ndim == 2 and qk > 0 and qk % 2 == 0 and (qk // 2) % K == 0","tryCatchPattern":null,"preventionTips":["Derive HV, V, K from initial_state.shape[-3:] and pass them consistently","Build mixed_qkv only via torch.cat([q, k, v], dim=1) from reshaped per-head tensors"],"tags":["shape-validation","triton-kernel","gated-delta-rule","packed-decode"],"backgroundTag":"tensor-shape-mismatch","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}