{"record":{"id":"9dc9da1ff5ab745f","repo":"sgl-project/sglang","slug":"dspark-kda-mtp-requires-a-fixed-1-num-spec-dense","errorCode":null,"errorMessage":"DSpARK KDA MTP requires a fixed 1 + num_spec dense tokens per request; got T={T}, N={N}","messagePattern":"DSpARK KDA MTP requires a fixed 1 \\+ num_spec dense tokens per request; got T=(.+?), N=(.+?)","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"python/sglang/kernels/ops/kimi_k3/kda_decode_mtp.py","lineNumber":997,"sourceCode":"    \"\"\"\n    import torch\n\n    H = x_q.shape[2]\n    N = cu_seqlens.numel() - 1\n    T = x_q.shape[1]\n    expected_shape = (1, T, H, TILE_K)\n    if tuple(x_q.shape) != expected_shape or tuple(x_k.shape) != expected_shape:\n        raise ValueError(f\"expected q/k shape {expected_shape}\")\n    if tuple(x_v.shape) != expected_shape or tuple(g.shape) != expected_shape:\n        raise ValueError(f\"expected v/g shape {expected_shape}\")\n    if tuple(beta.shape) != (1, T, H):\n        raise ValueError(f\"expected beta shape {(1, T, H)}\")\n    # T // N == 1 is num_spec == 0: one token per request, i.e. a plain decode\n    # step. The backend never dispatches here for it (that is the dedicated\n    # decode kernel's job), but the layout is legal and benchmarks compare the\n    # two at this point, so the wrapper accepts it.\n    if N <= 0 or T % N != 0 or T // N < 1:\n        raise ValueError(\n            f\"DSpARK KDA MTP requires a fixed 1 + num_spec dense tokens per \"\n            f\"request; got T={T}, N={N}\"\n        )\n    num_spec = T // N - 1\n    if recurrent_state.shape[1:] != (H, TILE_K, TILE_K):\n        raise ValueError(\"expected recurrent state layout [pool, H, V=128, K=128]\")\n    if (\n        recurrent_state.dtype != torch.float32\n        or tuple(recurrent_state.stride()[-3:]) != (TILE_K * TILE_K, TILE_K, 1)\n        or recurrent_state.stride(0) % 4 != 0\n        or recurrent_state.storage_offset() % 4 != 0\n    ):\n        raise ValueError(\n            \"cp.async recurrent state requires fp32 contiguous [H, V, K] \"\n            \"inner layout and 16-byte-aligned slot offsets\"\n        )\n    rings = (replayssm_rawv, replayssm_rawk, replayssm_g, replayssm_beta)\n    cache_ring = all(ring is not None for ring in rings)","sourceCodeStart":979,"sourceCodeEnd":1015,"githubUrl":"https://github.com/sgl-project/sglang/blob/0132848349585cfe6aae51c4941cbae872505f8a/python/sglang/kernels/ops/kimi_k3/kda_decode_mtp.py#L979-L1015","documentation":"The MTP variant of the kernel requires every request to carry exactly the same dense token count: T (flattened tokens) must be divisible by N (number of requests from cu_seqlens), giving 1 + num_spec tokens each. Violations (T % N != 0, N <= 0, or ragged requests) raise this error.","triggerScenarios":"Calling fused_kda_decode_mtp_dspark where cu_seqlens segments have unequal lengths (ragged speculative draft counts), N=0 (empty cu_seqlens), or T not a multiple of N.","commonSituations":"Mixed num_spec per request after dynamic speculative-length scheduling; a cu_seqlens built for plain decode concatenated with MTP drafts; off-by-one in cu_seqlens producing an extra empty segment.","solutions":["Ensure all requests have identical token counts (1 + num_spec) and T == N * (1 + num_spec)","Pad requests with fewer draft tokens or route ragged batches to a non-MTP KDA decode path","Validate cu_seqlens: monotonically increasing, N >= 1, differences all equal"],"exampleFix":"# before\nfused_kda_decode_mtp_dspark(q, k, v, g, beta, cu_seqlens, state)  # ragged lengths\n# after\nassert cu_seqlens.numel() >= 2 and T % (cu_seqlens.numel() - 1) == 0\n# or route ragged batches elsewhere:\nif T % N != 0:\n    return kda_decode_fallback(...)","handlingStrategy":"validation","validationCode":"N = cu_seqlens.numel() - 1\nassert N > 0 and T % N == 0 and T // N >= 1","typeGuard":"def uniform_mtp_batch(cu_seqlens: torch.Tensor) -> bool:\n    lens = cu_seqlens[1:] - cu_seqlens[:-1]\n    return lens.numel() > 0 and bool((lens == lens[0]).all())","tryCatchPattern":null,"preventionTips":["Pad requests to a uniform 1 + num_spec token count before calling","Route ragged batches to the generic KDA decode backend","Fuzz-test cu_seqlens builders for off-by-one empty segments"],"tags":["kda","mtp","cu-seqlens","uniform-batch"],"backgroundTag":"inconsistent-batch-layout","analyzedSha":"0132848349585cfe6aae51c4941cbae872505f8a","analyzedAt":"2026-08-28T05:10:05.995Z","schemaVersion":2},"datasetVersion":"2026-08-28T06:17:29.519Z"}