sgl-project/sglang · error · ValueError

Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V

Error message

Invalid packed `mixed_qkv` last dim={qkv_dim} for HV={HV}, V={V}.

What it means

Thrown by fused_recurrent_gated_delta_rule_packed_decode when the packed mixed_qkv tensor's last dimension cannot be split into Q/K and V parts. The kernel computes qk_dim = mixed_qkv.shape[1] - HV*V and requires it to be positive and even (Q and K have equal head dim). If the layout or head counts don't match the expected packed q|k|v format, this validation fails.

Source

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

        raise ValueError("`initial_state` must be contiguous in the last dim.")
    HV, V, K = initial_state.shape[-3:]
    if a.shape[1] != HV or b.shape[1] != HV:
        raise ValueError(
            f"`a`/`b` must have shape [B, HV] with HV={HV} (got a.shape={tuple(a.shape)}, b.shape={tuple(b.shape)})."
        )
    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

View on GitHub (pinned to 0132848349)

Solutions

1. Check that mixed_qkv.shape[1] == HV*V + 2*H*K for your model config and that qk_dim = shape[1] - HV*V is even and > 0
2. Verify the V, HV, K values inferred from initial_state.shape[-3:] match the model's projection layer sizes
3. Ensure mixed_qkv was built by concatenating q, k, v along dim 1 with no extra columns

Example fix

# before
mixed_qkv = torch.cat([q, k, v, gate], dim=1)  # extra gate column
# after
mixed_qkv = torch.cat([q, k, v], dim=1)  # qk_dim = 2*H*K, even

qkv_dim = mixed_qkv.shape[1]
qk_dim = qkv_dim - HV * V
assert qk_dim > 0 and qk_dim % 2 == 0, (qkv_dim, HV, V)

def valid_packed_qkv(mixed_qkv, HV, V, K):
    qk = mixed_qkv.shape[1] - HV * V
    return mixed_qkv.ndim == 2 and qk > 0 and qk % 2 == 0 and (qk // 2) % K == 0

When it happens

Trigger: Calling fused_recurrent_gated_delta_rule_packed_decode with a mixed_qkv tensor whose shape[1] is not HV*V + 2*H*K, e.g. wrong V/HV values, an extra gate column packed in, or a transposed/non-packed tensor.

Common situations: Model config mismatch (V, HV, K from initial_state not matching the projection output width), changing head_dim or num_v_heads without regenerating the packed qkv buffer, or packing additional tensors (e.g. gates) into mixed_qkv.

Related errors

• Invalid packed Q size {q_dim}: must be divisible by K={K}.
• 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/3f079a5dfcf31faa. Report an issue: GitHub.