{"record":{"id":"d40ac57efaaab279","repo":"sgl-project/sglang","slug":"invalid-head-config-inferred-from-mixed-qkv-h-h","errorCode":null,"errorMessage":"Invalid head config inferred from mixed_qkv: H={H}, HV={HV}.","messagePattern":"Invalid head config inferred from mixed_qkv: H=(.+?), HV=(.+?)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/fla/fused_recurrent.py","lineNumber":352,"sourceCode":"            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)\n    stride_init_state_token = initial_state.stride(0)\n    stride_final_state_token = initial_state.stride(0)\n    stride_indices_seq = ssm_state_indices.stride(0)","sourceCodeStart":334,"sourceCodeEnd":370,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/fla/fused_recurrent.py#L334-L370","documentation":"The head count H inferred as q_dim//K must be positive and must evenly divide HV (the number of value heads), since the kernel maps each value head to a group of query heads (GQA-style). If HV % H != 0, the query-to-value head grouping is invalid.","triggerScenarios":"Calling packed decode with a config where num_q_heads does not divide num_v_heads (e.g. H=7, HV=4), or where inferred H<=0 due to a tiny mixed_qkv.","commonSituations":"Custom model configs with non-GQA head ratios, or a truncated mixed_qkv tensor making q_dim smaller than K.","solutions":["Ensure num_query_heads divides num_v_heads (standard GQA ratio)","Check mixed_qkv.shape[1] is large enough that q_dim//K >= 1","Align the model's num_attention_heads / num_key_value_heads with the packed layout"],"exampleFix":null,"handlingStrategy":"validation","validationCode":"H = ((mixed_qkv.shape[1] - HV * V) // 2) // K\nassert H > 0 and HV % H == 0, (H, HV)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Ensure num_q_heads divides num_v_heads in the model config","Validate head config once at model init, not per decode step"],"tags":["shape-validation","gqa","triton-kernel","packed-decode"],"backgroundTag":"gqa-head-ratio-invalid","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}