Introduce global alloc_len_per_decode & clean check decode memory (#15115)
This commit is contained in:
@@ -1,7 +1,7 @@
|
||||
import logging
|
||||
from copy import copy
|
||||
from dataclasses import dataclass
|
||||
from typing import ClassVar, List, Optional, Tuple
|
||||
from typing import List, Optional, Tuple
|
||||
|
||||
import torch
|
||||
import torch.nn.functional as F
|
||||
@@ -614,9 +614,6 @@ class EagleVerifyInput(SpecInput, EagleVerifyInputV2Mixin):
|
||||
|
||||
@dataclass
|
||||
class EagleDraftInput(SpecInput, EagleDraftInputV2Mixin):
|
||||
# Constant: alloc length per decode step
|
||||
ALLOC_LEN_PER_DECODE: ClassVar[int] = None
|
||||
|
||||
# The inputs for decode
|
||||
# shape: (b, topk)
|
||||
topk_p: torch.Tensor = None
|
||||
|
||||
@@ -10,6 +10,7 @@ import triton.language as tl
|
||||
|
||||
from sglang.srt.layers.logits_processor import LogitsProcessorOutput
|
||||
from sglang.srt.managers.schedule_batch import ModelWorkerBatch, ScheduleBatch
|
||||
from sglang.srt.managers.utils import get_alloc_len_per_decode
|
||||
from sglang.srt.mem_cache.common import (
|
||||
alloc_paged_token_slots_extend,
|
||||
alloc_token_slots,
|
||||
@@ -92,9 +93,10 @@ class EagleDraftInputV2Mixin:
|
||||
cur_kv_lens_cpu = []
|
||||
nxt_kv_lens_cpu = []
|
||||
num_needed_tokens = 0
|
||||
alloc_len_per_decode = get_alloc_len_per_decode()
|
||||
for r in batch.reqs:
|
||||
# Over-allocation happens here
|
||||
x = r.kv_committed_len + 2 * self.ALLOC_LEN_PER_DECODE - r.kv_allocated_len
|
||||
x = r.kv_committed_len + 2 * alloc_len_per_decode - r.kv_allocated_len
|
||||
cur_kv_lens_cpu.append(r.kv_allocated_len)
|
||||
nxt_kv_lens_cpu.append(r.kv_allocated_len + x)
|
||||
num_needed_tokens += x
|
||||
|
||||
@@ -394,10 +394,11 @@ class EAGLEWorker(TpModelWorker):
|
||||
# [ topk 0 ] [ topk 1 ]
|
||||
# [iter=0, iter=1, iter=2] [iter=0, iter=1, iter=2]
|
||||
if self.page_size == 1:
|
||||
alloc_len_per_decode = self.speculative_num_steps * self.topk
|
||||
# TODO: We only need self.speculative_num_steps - 1 * topk cache loc
|
||||
out_cache_loc, token_to_kv_pool_state_backup = alloc_token_slots(
|
||||
batch.tree_cache,
|
||||
num_seqs * self.speculative_num_steps * self.topk,
|
||||
num_seqs * alloc_len_per_decode,
|
||||
backup_state=True,
|
||||
)
|
||||
else:
|
||||
|
||||
@@ -107,11 +107,6 @@ class EagleDraftWorker(BaseDraftWorker):
|
||||
server_args.speculative_algorithm
|
||||
)
|
||||
|
||||
# Set constant
|
||||
EagleDraftInput.ALLOC_LEN_PER_DECODE = max(
|
||||
self.speculative_num_steps * self.topk, self.speculative_num_draft_tokens
|
||||
)
|
||||
|
||||
# Do not capture cuda graph in `TpModelWorker` init,
|
||||
# will capture later with init_cuda_graphs()
|
||||
backup_disable_cuda_graph = server_args.disable_cuda_graph
|
||||
|
||||
Reference in New Issue
Block a user