sgl-project/sglang · error · ValueError

Packed decode kernel only supports NK=1 (got K={K}, BK={BK})

Error message

Packed decode kernel only supports NK=1 (got K={K}, BK={BK}).

What it means

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.

Source

Thrown at python/sglang/kernels/ops/attention/fla/fused_recurrent.py:358

    qkv_dim = mixed_qkv.shape[1]
    qk_dim = qkv_dim - HV * V
    if qk_dim <= 0 or qk_dim % 2 != 0:
        raise ValueError(
            f"Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V={V}."
        )
    q_dim = qk_dim // 2
    if q_dim % K != 0:
        raise ValueError(f"Invalid packed Q size {q_dim}: must be divisible by K={K}.")
    H = q_dim // K
    if H <= 0 or HV % H != 0:
        raise ValueError(
            f"Invalid head config inferred from mixed_qkv: H={H}, HV={HV}."
        )

    BK = triton.next_power_of_2(K)
    if triton.cdiv(K, BK) != 1:
        raise ValueError(
            f"Packed decode kernel only supports NK=1 (got K={K}, BK={BK})."
        )
    BV = min(triton.next_power_of_2(V), 32)
    num_stages = 3
    num_warps = 1

    stride_mixed_qkv_tok = mixed_qkv.stride(0)
    stride_a_tok = a.stride(0)
    stride_b_tok = b.stride(0)
    stride_init_state_token = initial_state.stride(0)
    stride_final_state_token = initial_state.stride(0)
    stride_indices_seq = ssm_state_indices.stride(0)

    NV = triton.cdiv(V, BV)
    grid = (NV, B * HV)
    fused_recurrent_gated_delta_rule_packed_decode_kernel[grid](
        mixed_qkv=mixed_qkv,
        a=a,

View on GitHub (pinned to 0132848349)

Solutions

1. Check initial_state.shape[-1] (K) is a sane positive head dim (e.g. 64-256)
2. If head_dim is very large, use the non-packed/unfused decode path instead
3. Avoid patching BK to a smaller cap than K

BK = triton.next_power_of_2(K)
assert triton.cdiv(K, BK) == 1, (K, BK)

try:
    fused_recurrent_gated_delta_rule_packed_decode(...)
except ValueError:
    out = unfused_decode_step(...)

When it happens

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

Common situations: Degenerate initial_state with K=0, or a patched kernel that clamps BK below K (e.g. very large head_dim > 256).

Related errors

• Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V
• Invalid packed Q size {q_dim}: must be divisible by K={K}.
• Invalid head config inferred from mixed_qkv: H={H}, HV={HV}.

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