sgl-project/sglang · error · NotImplementedError

eqlen with B>1 and T % {BT} != 0 not supported (got B={B}, T

Error message

eqlen with B>1 and T % {BT} != 0 not supported (got B={B}, T={T}).

What it means

chunk_kda_fwd pads non-varlen (eqlen) sequences whose total token count T is not a multiple of _K123_CHUNKS_PER_BLOCK*BT. That padding path only works when either B==1 or T is already a multiple of BT, because multi-seq eqlen batches cannot be host-padded without repacking memory. Any eqlen batch with B>1 and T % BT != 0 is rejected before launch.

Source

Thrown at python/sglang/kernels/ops/attention/linear/kda_nvidia_prefill/chunk_fwd.py:924

    V_dim = v.shape[-1]
    device = q.device
    BT = 64

    # Phase 1: the eqlen persistent scheduler assigns CHUNKS_PER_BLOCK chunks
    # to each workgroup, so pad T to that scheduling unit even when T is
    # already a 64-row chunk multiple. This prevents zero-workgroup launches
    # for short inputs and avoids silently dropping a trailing chunk group.
    # K123 eqlen runs unchanged and sees only full 64-row chunks; zero/sentinel
    # padding handles the caller's boundary at the data level.
    #
    # Varlen with non-aligned seq lengths is a separate problem (Phase 2):
    # multi-seq varlen can't be host-padded without repacking memory.
    real_T = T
    CPB_BT = _K123_CHUNKS_PER_BLOCK * BT
    needs_eqlen_pad = (not is_varlen) and (T % CPB_BT != 0)
    if needs_eqlen_pad:
        if B != 1 and T % BT != 0:
            raise NotImplementedError(
                f"eqlen with B>1 and T % {BT} != 0 not supported "
                f"(got B={B}, T={T})."
            )
        T_padded = ((T + CPB_BT - 1) // CPB_BT) * CPB_BT
        # Pre-allocated padded scratch buffers (per (B,T_padded,H,K,dtype) cache
        # key). torch.cat would reallocate + copy the full 200MB q tensor every
        # call — caching the destination buffer drops that to a single slice
        # copy of the valid prefix (caller already lives in our buffer for
        # subsequent calls reusing the same id, but we re-copy unconditionally
        # since the caller may have updated the data in-place).
        q_pad, k_pad, v_pad, g_pad, beta_pad = _get_padded_input_buffers(
            B, T_padded, H, K, q.dtype, g.dtype, beta.dtype, q.device, real_T
        )
        # q/k/v/beta zero-padded → K1/K2 MMAs naturally produce 0 for OOB rows.
        # Tail [real_T:] of q_pad/k_pad/v_pad/beta_pad is pre-zeroed at cache
        # init and never written, so we only copy the valid prefix.
        q_pad[:, :real_T].copy_(q)
        k_pad[:, :real_T].copy_(k)

View on GitHub (pinned to 0132848349)

Solutions

1. Switch to varlen inputs (pass cu_seqlens/is_varlen path) so padding is not needed
2. Reshape/pad the batch so T % BT == 0 before calling, or run with B == 1
3. Repack into a B=1 packed layout (concatenate sequences and use varlen offsets)

Example fix

// before
out = chunk_kda_fwd(q, k, v, ..., B=4, T=1000)  # 1000 % 64 != 0
// after
# pad T up to a multiple of BT (or use the varlen path)
T_pad = ((T + BT - 1) // BT) * BT
q = F.pad(q, (0, 0, 0, 0, 0, T_pad - T))
out = chunk_kda_fwd(q, k, v, ..., B=4, T=T_pad)

BT = get_kda_block_tile(cfg)  # same BT the kernel uses
if not is_varlen and B > 1 and T % BT != 0:
    T = ((T + BT - 1) // BT) * BT  # pad, or repack as varlen/B=1

catch NotImplementedError from chunk_kda_fwd and fall back to the varlen/packed layout

When it happens

Trigger: Calling chunk_kda_fwd (chunk_fwd.py) with is_varlen=False, batch size B>1, and a sequence length T that is not divisible by the block tile BT (BT depends on kernel config, e.g. 64/128).

Common situations: Serving mixed-length non-varlen batches of KDA models where the summed/major T happens to be an odd multiple; benchmark scripts that build eqlen inputs of arbitrary lengths; changing chunk/block-tile config so T no longer aligns.

Related errors

• KDA cutedsl: safe_gate (lower_bound) not yet supported
• CuteDSLKDAKernel does not support target_verify
• Unexpected A_log shape: {A_log.shape}; expected numel={HV}
• Unexpected dt_bias shape: {dt_bias.shape}; expected numel={H
• `dt_bias` must have {HV * K} elements (got {dt_bias.numel()}

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