{"record":{"id":"ca55de74bfb1bad0","repo":"sgl-project/sglang","slug":"invalid-packed-q-size-q-dim-must-be-divisible-b","errorCode":null,"errorMessage":"Invalid packed Q size {q_dim}: must be divisible by K={K}.","messagePattern":"Invalid packed Q size (.+?): must be divisible by K=(.+?)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/fla/fused_recurrent.py","lineNumber":349,"sourceCode":"        )\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\n    num_warps = 1\n\n    stride_mixed_qkv_tok = mixed_qkv.stride(0)\n    stride_a_tok = a.stride(0)\n    stride_b_tok = b.stride(0)","sourceCodeStart":331,"sourceCodeEnd":367,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/fla/fused_recurrent.py#L331-L367","documentation":"After splitting mixed_qkv, the inferred Q dimension (qk_dim//2) must be a multiple of the key head dim K. This error means q and k head dims are inconsistent: the total Q width is not divisible by per-head K, so head count H cannot be computed.","triggerScenarios":"qk_dim/2 is not divisible by K, e.g. K=128 but total Q width is 96, or mixed_qkv includes asymmetric q/k sizes.","commonSituations":"Config typos where key_head_dim differs from query_head_dim packing assumptions, or partial/offset packing of q into mixed_qkv.","solutions":["Verify K (from initial_state last dim) equals the model's per-head key dim used to pack q","Ensure q occupies exactly H*K columns and k occupies H*K columns in mixed_qkv","Re-pack mixed_qkv as cat([q.reshape(B,-1), k.reshape(B,-1), v.reshape(B,-1)], dim=1)"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"q_dim = (mixed_qkv.shape[1] - HV * V) // 2\nassert q_dim % K == 0, (q_dim, K)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Keep query and key head dims equal when packing mixed_qkv","Unit-test the packing helper against model config (H, K, HV, V)"],"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"}