sgl-project/sglang · error · ValueError

Invalid packed Q size {q_dim}: must be divisible by K={K}.

Error message

Invalid packed Q size {q_dim}: must be divisible by K={K}.

What it means

After splitting mixed_qkv, the inferred Q dimension (qk_dim//2) must be a multiple of the key head dim K. This error means q and k head dims are inconsistent: the total Q width is not divisible by per-head K, so head count H cannot be computed.

Source

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

        )
    if A_log.numel() != HV or dt_bias.numel() != HV:
        raise ValueError(
            f"`A_log` and `dt_bias` must have {HV} elements (got A_log.numel()={A_log.numel()}, dt_bias.numel()={dt_bias.numel()})."
        )
    if out.shape != (B, 1, HV, V):
        raise ValueError(
            f"`out` must have shape {(B, 1, HV, V)} (got out.shape={tuple(out.shape)})."
        )

    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)

View on GitHub (pinned to 0132848349)

Solutions

1. Verify K (from initial_state last dim) equals the model's per-head key dim used to pack q
2. Ensure q occupies exactly H*K columns and k occupies H*K columns in mixed_qkv
3. Re-pack mixed_qkv as cat([q.reshape(B,-1), k.reshape(B,-1), v.reshape(B,-1)], dim=1)

q_dim = (mixed_qkv.shape[1] - HV * V) // 2
assert q_dim % K == 0, (q_dim, K)

When it happens

Trigger: qk_dim/2 is not divisible by K, e.g. K=128 but total Q width is 96, or mixed_qkv includes asymmetric q/k sizes.

Common situations: Config typos where key_head_dim differs from query_head_dim packing assumptions, or partial/offset packing of q into mixed_qkv.

Related errors

• Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V
• Invalid head config inferred from mixed_qkv: H={H}, HV={HV}.
• The number of initial states is expected to be equal to the
• Packed decode kernel only supports NK=1 (got K={K}, BK={BK})
• The pointers must be multiple of 16 bytes.

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