sgl-project/sglang · error · ValueError

Unsupported split_k: {split_k}

Error message

Unsupported split_k: {split_k}

What it means

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.

Source

Thrown at python/sglang/kernels/ops/attention/nsa_triton_decode/triton_mla_kernels_decode_fused.py:742

                output.stride(0),
                output.stride(1),
                output.stride(2),
                lse.stride(0),
                lse.stride(1),
                HAS_ATTN_SINK=attn_sink is not None,
            )
        else:
            BLOCK_H_COMBINE = 16
            BLOCK_D_COMBINE = 128
            grid_combine = (total_tokens, triton.cdiv(h_q, BLOCK_H_COMBINE))

            # Select appropriate combine kernel based on split_k
            if split_k == 2:
                combine_kernel = _combine_splitk_kernel_2
            elif split_k == 4:
                combine_kernel = _combine_splitk_kernel
            else:
                raise ValueError(f"Unsupported split_k: {split_k}")

            combine_kernel[grid_combine](
                partial_output,
                partial_lse,
                attn_sink_tensor,
                output,
                lse,
                total_tokens,
                _bucket_total_tokens(total_tokens),
                h_q,
                d_v,
                partial_output.stride(0),
                partial_output.stride(1),
                partial_output.stride(2),
                partial_output.stride(3),
                partial_lse.stride(0),
                partial_lse.stride(1),
                partial_lse.stride(2),

View on GitHub (pinned to 0132848349)

Solutions

1. Set split_k to 2 or 4
2. If split_k == 1 is desired, use the non-split code path / kernel variant that skips the combine stage

Example fix

# before
out = fused_gather_attn_decode_dsv4(..., split_k=3)
# after
out = fused_gather_attn_decode_dsv4(..., split_k=4)

assert split_k in (2, 4), f'split_k must be 2 or 4, got {split_k}'
out = fused_gather_attn_decode_dsv4(..., split_k=split_k)

When it happens

Trigger: 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.

Common situations: 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.

Related errors

• Unsupported d_qk: {d_qk}. Expected {DSV4_D_QK} (DeepSeek V4)
• Unsupported content type ${header.content_type}
• Unknown serve backend {name!r}. Available values: {available
• k_cache can only be None when only_qv=True
• q can only be None when only_qv=True

AI-assisted analysis of sgl-project/sglang@0132848349 (2026-08-28). Data as JSON: /api/errors/88481d0ac5045e56. Report an issue: GitHub.