{"record":{"id":"5d0ca3e3d3534178","repo":"sgl-project/sglang","slug":"packed-decode-kernel-only-supports-nk-1-got-k-k","errorCode":null,"errorMessage":"Packed decode kernel only supports NK=1 (got K={K}, BK={BK}).","messagePattern":"Packed decode kernel only supports NK=1 \\(got K=(.+?), BK=(.+?)\\)\\.","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/fla/fused_recurrent.py","lineNumber":358,"sourceCode":"\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)\n\n    NV = triton.cdiv(V, BV)\n    grid = (NV, B * HV)\n    fused_recurrent_gated_delta_rule_packed_decode_kernel[grid](\n        mixed_qkv=mixed_qkv,\n        a=a,","sourceCodeStart":340,"sourceCodeEnd":376,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/fla/fused_recurrent.py#L340-L376","documentation":"The packed decode Triton kernel processes the full K dimension in a single block, so it requires K <= BK = next_power_of_2(K), i.e. cdiv(K,BK)==1. This always holds mathematically for next_power_of_2 unless K exceeds the block-size limit or is 0/invalid, so hitting this error indicates an anomalous K.","triggerScenarios":"K with an invalid value (0 or degenerate) where next_power_of_2 rounding makes cdiv(K,BK) != 1, or a modified kernel with a BK cap.","commonSituations":"Degenerate initial_state with K=0, or a patched kernel that clamps BK below K (e.g. very large head_dim > 256).","solutions":["Check initial_state.shape[-1] (K) is a sane positive head dim (e.g. 64-256)","If head_dim is very large, use the non-packed/unfused decode path instead","Avoid patching BK to a smaller cap than K"],"exampleFix":null,"handlingStrategy":"fallback","validationCode":"BK = triton.next_power_of_2(K)\nassert triton.cdiv(K, BK) == 1, (K, BK)","typeGuard":null,"tryCatchPattern":"try:\n    fused_recurrent_gated_delta_rule_packed_decode(...)\nexcept ValueError:\n    out = unfused_decode_step(...)","preventionTips":["Keep head_dim K within typical 64-256 range","Fall back to the unfused reference path for exotic head dims"],"tags":["kernel-limit","triton-kernel","packed-decode"],"backgroundTag":"kernel-block-size-limit","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}