{"record":{"id":"88481d0ac5045e56","repo":"sgl-project/sglang","slug":"unsupported-split-k-split-k","errorCode":null,"errorMessage":"Unsupported split_k: {split_k}","messagePattern":"Unsupported split_k: (.+?)","errorType":"exception","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/nsa_triton_decode/triton_mla_kernels_decode_fused.py","lineNumber":742,"sourceCode":"                output.stride(0),\n                output.stride(1),\n                output.stride(2),\n                lse.stride(0),\n                lse.stride(1),\n                HAS_ATTN_SINK=attn_sink is not None,\n            )\n        else:\n            BLOCK_H_COMBINE = 16\n            BLOCK_D_COMBINE = 128\n            grid_combine = (total_tokens, triton.cdiv(h_q, BLOCK_H_COMBINE))\n\n            # Select appropriate combine kernel based on split_k\n            if split_k == 2:\n                combine_kernel = _combine_splitk_kernel_2\n            elif split_k == 4:\n                combine_kernel = _combine_splitk_kernel\n            else:\n                raise ValueError(f\"Unsupported split_k: {split_k}\")\n\n            combine_kernel[grid_combine](\n                partial_output,\n                partial_lse,\n                attn_sink_tensor,\n                output,\n                lse,\n                total_tokens,\n                _bucket_total_tokens(total_tokens),\n                h_q,\n                d_v,\n                partial_output.stride(0),\n                partial_output.stride(1),\n                partial_output.stride(2),\n                partial_output.stride(3),\n                partial_lse.stride(0),\n                partial_lse.stride(1),\n                partial_lse.stride(2),","sourceCodeStart":724,"sourceCodeEnd":760,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/nsa_triton_decode/triton_mla_kernels_decode_fused.py#L724-L760","documentation":"In fused_gather_attn_decode_dsv4, after a split-K decode the partial results must be combined by one of two specialized combine kernels that only exist for split_k == 2 or split_k == 4. Any other split_k value is rejected with ValueError.","triggerScenarios":"Calling fused_gather_attn_decode_dsv4 with a split_k argument other than 2 or 4 (e.g. 1, 3, 8) that reaches the combine-kernel selection branch.","commonSituations":"Tuning split_k for short-sequence NSA/MLA decode and trying values the combine kernels were not built for; copying configs between the dsv4 kernel variants that accept different split_k sets.","solutions":["Set split_k to 2 or 4","If split_k == 1 is desired, use the non-split code path / kernel variant that skips the combine stage"],"exampleFix":"# before\nout = fused_gather_attn_decode_dsv4(..., split_k=3)\n# after\nout = fused_gather_attn_decode_dsv4(..., split_k=4)","handlingStrategy":"validation","validationCode":"assert split_k in (2, 4), f'split_k must be 2 or 4, got {split_k}'\nout = fused_gather_attn_decode_dsv4(..., split_k=split_k)","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Whitelist split_k values from tuner output","Share a single validated split_k constant across the dsv4 decode call sites"],"tags":["nsa","triton","split-k","invalid-argument"],"backgroundTag":"unsupported-parameter-value","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}