{"record":{"id":"0be76f3668833950","repo":"sgl-project/sglang","slug":"eqlen-with-b-1-and-t-bt-0-not-supported-go","errorCode":null,"errorMessage":"eqlen with B>1 and T % {BT} != 0 not supported (got B={B}, T={T}).","messagePattern":"eqlen with B>1 and T % (.+?) != 0 not supported \\(got B=(.+?), T=(.+?)\\)\\.","errorType":"exception","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/attention/linear/kda_nvidia_prefill/chunk_fwd.py","lineNumber":924,"sourceCode":"    V_dim = v.shape[-1]\n    device = q.device\n    BT = 64\n\n    # Phase 1: the eqlen persistent scheduler assigns CHUNKS_PER_BLOCK chunks\n    # to each workgroup, so pad T to that scheduling unit even when T is\n    # already a 64-row chunk multiple. This prevents zero-workgroup launches\n    # for short inputs and avoids silently dropping a trailing chunk group.\n    # K123 eqlen runs unchanged and sees only full 64-row chunks; zero/sentinel\n    # padding handles the caller's boundary at the data level.\n    #\n    # Varlen with non-aligned seq lengths is a separate problem (Phase 2):\n    # multi-seq varlen can't be host-padded without repacking memory.\n    real_T = T\n    CPB_BT = _K123_CHUNKS_PER_BLOCK * BT\n    needs_eqlen_pad = (not is_varlen) and (T % CPB_BT != 0)\n    if needs_eqlen_pad:\n        if B != 1 and T % BT != 0:\n            raise NotImplementedError(\n                f\"eqlen with B>1 and T % {BT} != 0 not supported \"\n                f\"(got B={B}, T={T}).\"\n            )\n        T_padded = ((T + CPB_BT - 1) // CPB_BT) * CPB_BT\n        # Pre-allocated padded scratch buffers (per (B,T_padded,H,K,dtype) cache\n        # key). torch.cat would reallocate + copy the full 200MB q tensor every\n        # call — caching the destination buffer drops that to a single slice\n        # copy of the valid prefix (caller already lives in our buffer for\n        # subsequent calls reusing the same id, but we re-copy unconditionally\n        # since the caller may have updated the data in-place).\n        q_pad, k_pad, v_pad, g_pad, beta_pad = _get_padded_input_buffers(\n            B, T_padded, H, K, q.dtype, g.dtype, beta.dtype, q.device, real_T\n        )\n        # q/k/v/beta zero-padded → K1/K2 MMAs naturally produce 0 for OOB rows.\n        # Tail [real_T:] of q_pad/k_pad/v_pad/beta_pad is pre-zeroed at cache\n        # init and never written, so we only copy the valid prefix.\n        q_pad[:, :real_T].copy_(q)\n        k_pad[:, :real_T].copy_(k)","sourceCodeStart":906,"sourceCodeEnd":942,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/attention/linear/kda_nvidia_prefill/chunk_fwd.py#L906-L942","documentation":"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.","triggerScenarios":"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).","commonSituations":"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.","solutions":["Switch to varlen inputs (pass cu_seqlens/is_varlen path) so padding is not needed","Reshape/pad the batch so T % BT == 0 before calling, or run with B == 1","Repack into a B=1 packed layout (concatenate sequences and use varlen offsets)"],"exampleFix":"// before\nout = chunk_kda_fwd(q, k, v, ..., B=4, T=1000)  # 1000 % 64 != 0\n// after\n# pad T up to a multiple of BT (or use the varlen path)\nT_pad = ((T + BT - 1) // BT) * BT\nq = F.pad(q, (0, 0, 0, 0, 0, T_pad - T))\nout = chunk_kda_fwd(q, k, v, ..., B=4, T=T_pad)","handlingStrategy":"validation","validationCode":"BT = get_kda_block_tile(cfg)  # same BT the kernel uses\nif not is_varlen and B > 1 and T % BT != 0:\n    T = ((T + BT - 1) // BT) * BT  # pad, or repack as varlen/B=1","typeGuard":null,"tryCatchPattern":"catch NotImplementedError from chunk_kda_fwd and fall back to the varlen/packed layout","preventionTips":["Prefer varlen (cu_seqlens) inputs for mixed-length batches","Keep eqlen batch T aligned to the kernel block tile","Never assume arbitrary (B, T) combos are supported in eqlen mode"],"tags":["kda","linear-attention","shape-mismatch","not-implemented"],"backgroundTag":"sequence-length-not-divisible-by-block","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}