feat: add req param cache for decode

Author: niushengxiao

lightllm/common/basemodel/triton_kernel/apply_penalty.py

@@ -19,9 +19,9 @@ def _fwd_kernel_apply_penalty(
     eos_ids,
     mask_eos_reqs,
     stride_logit_b,
-    stride_logit_s,
     BLOCK_P: tl.constexpr,
     EOS_ID_NUM: tl.constexpr,
+    IS_EOS_PENALTY: tl.constexpr,
 ):
     cur_batch = tl.program_id(0)
     cur_freqency = tl.load(freqency_penalty + cur_batch)
@@ -46,18 +46,19 @@ def _fwd_kernel_apply_penalty(
         output_ptr = Logits + cur_batch * stride_logit_b + batch_ids
         tl.store(output_ptr, pre_logits, mask=cur_batch_id_offset < cur_batch_end_index)
 
-    mask_eos = tl.load(mask_eos_reqs + cur_batch)
-    exponential_decay_length_penalty = tl.load(exponential_decay_length_penalties + cur_batch)
-    length_penalty = tl.load(length_penalty_idx + cur_batch)
-    penalty_scale = tl.exp2(tl.log2(exponential_decay_length_penalty) * length_penalty) - 1
+    if IS_EOS_PENALTY:
+        mask_eos = tl.load(mask_eos_reqs + cur_batch)
+        exponential_decay_length_penalty = tl.load(exponential_decay_length_penalties + cur_batch)
+        length_penalty = tl.load(length_penalty_idx + cur_batch)
+        penalty_scale = tl.exp2(tl.log2(exponential_decay_length_penalty) * length_penalty) - 1
 
-    for eos_index in range(EOS_ID_NUM):
-        eos_id = tl.load(eos_ids + eos_index)
-        cur_eos_logit_ptr = Logits + cur_batch * stride_logit_b + eos_id
-        cur_eos_logit = tl.load(cur_eos_logit_ptr)
-        cur_eos_logit = cur_eos_logit + tl.abs(cur_eos_logit) * penalty_scale
-        cur_eos_logit = tl.where(mask_eos, -10000000.0, cur_eos_logit)
-        tl.store(cur_eos_logit_ptr, cur_eos_logit)
+        for eos_index in range(EOS_ID_NUM):
+            eos_id = tl.load(eos_ids + eos_index)
+            cur_eos_logit_ptr = Logits + cur_batch * stride_logit_b + eos_id
+            cur_eos_logit = tl.load(cur_eos_logit_ptr)
+            cur_eos_logit = cur_eos_logit + tl.abs(cur_eos_logit) * penalty_scale
+            cur_eos_logit = tl.where(mask_eos, -10000000.0, cur_eos_logit)
+            tl.store(cur_eos_logit_ptr, cur_eos_logit)
     return
 
 
@@ -74,6 +75,7 @@ def apply_penalty(
     length_penalty_idx,
     eos_ids,
     mask_eos_reqs,
+    is_eos_penalty=False,
 ):
     assert Logits.is_contiguous()
     BLOCK_P = 1024
@@ -91,9 +93,9 @@ def apply_penalty(
         eos_ids,
         mask_eos_reqs,
         Logits.stride(0),
-        Logits.stride(1),
         num_warps=num_warps,
         BLOCK_P=BLOCK_P,
         EOS_ID_NUM=eos_ids.shape[0],
+        IS_EOS_PENALTY=is_eos_penalty,
     )
     return

lightllm/common/basemodel/triton_kernel/apply_penalty_cache.py

@@ -0,0 +1,198 @@
+import torch
+
+import triton
+import triton.language as tl
+import torch.nn.functional as F
+import numpy as np
+
+
+@triton.jit
+def _fwd_kernel_apply_penalty_cache(
+    Logits,
+    req_idxs,
+    presence_penalty,
+    freqency_penalty,
+    repetition_penalty,
+    p_token_vocabs,
+    stride_logit_b,
+    stride_p_token_vocabs_b,
+    BLOCK_P: tl.constexpr,
+):
+    cur_batch = tl.program_id(0)
+    block_idx = tl.program_id(1)
+    token_idx = tl.load(req_idxs + cur_batch)
+    cur_freqency = tl.load(freqency_penalty + token_idx)
+    cur_presence = tl.load(presence_penalty + token_idx)
+    cur_repetition = tl.load(repetition_penalty + token_idx)
+
+    batch_ids = BLOCK_P * block_idx + tl.arange(0, BLOCK_P)
+    batch_ids_count = tl.load(
+        p_token_vocabs + token_idx * stride_p_token_vocabs_b + batch_ids,
+        mask=batch_ids < stride_p_token_vocabs_b,
+        other=0,
+    )
+    row_start_ptr = Logits + cur_batch * stride_logit_b
+    cur_offset = row_start_ptr + batch_ids
+    cur_logits = tl.load(cur_offset, mask=batch_ids_count > 0, other=0.0)
+    rep_logits = tl.where(cur_logits > 0, cur_logits / cur_repetition, cur_logits * cur_repetition)
+    freq_logits = rep_logits - batch_ids_count * cur_freqency
+    pre_logits = freq_logits - cur_presence
+    output_ptr = Logits + cur_batch * stride_logit_b + batch_ids
+    tl.store(output_ptr, pre_logits, mask=batch_ids_count > 0)
+    return
+
+
+@triton.jit
+def _eos_penalty(
+    Logits,
+    req_idxs,
+    p_token_lens,
+    exponential_decay_length_penalties,
+    length_penalty_idx,
+    eos_ids,
+    mask_eos_reqs,
+    stride_logit_b,
+    EOS_ID_NUM: tl.constexpr,
+):
+    cur_batch = tl.program_id(0)
+    token_idx = tl.load(req_idxs + cur_batch)
+    exponential_decay_length_penalty = tl.load(exponential_decay_length_penalties + token_idx)
+    token_lens = tl.load(p_token_lens + cur_batch)
+    length_penalty = tl.maximum(token_lens - tl.load(length_penalty_idx + token_idx), 0)
+    penalty_scale = tl.exp2(tl.log2(exponential_decay_length_penalty) * length_penalty) - 1
+    mask_eos = tl.load(mask_eos_reqs + token_idx)
+    for eos_index in range(EOS_ID_NUM):
+        eos_id = tl.load(eos_ids + eos_index)
+        cur_eos_logit_ptr = Logits + cur_batch * stride_logit_b + eos_id
+        cur_eos_logit = tl.load(cur_eos_logit_ptr)
+        cur_eos_logit = cur_eos_logit + tl.abs(cur_eos_logit) * penalty_scale
+        cur_eos_logit = tl.where(token_lens < mask_eos, -10000000.0, cur_eos_logit)
+        tl.store(cur_eos_logit_ptr, cur_eos_logit)
+    return
+
+
+@torch.no_grad()
+def apply_penalty_cache(
+    Logits,
+    req_idxs,
+    presence_penalty,
+    freqency_penalty,
+    repetition_penalty,
+    p_token_vocabs,
+    p_token_lens,
+    exponential_decay_length_penalties,
+    length_penalty_idx,
+    eos_ids,
+    mask_eos_reqs,
+    is_eos_penalty=False,
+):
+    assert Logits.is_contiguous()
+    BLOCK_P = 1024
+    num_warps = 8
+    vocab_size = p_token_vocabs.shape[1]
+    _fwd_kernel_apply_penalty_cache[(Logits.shape[0], triton.cdiv(vocab_size, BLOCK_P))](
+        Logits,
+        req_idxs,
+        presence_penalty,
+        freqency_penalty,
+        repetition_penalty,
+        p_token_vocabs,
+        Logits.stride(0),
+        p_token_vocabs.stride(0),
+        num_warps=num_warps,
+        BLOCK_P=BLOCK_P,
+    )
+    if is_eos_penalty:
+        p_token_lens = p_token_vocabs[req_idxs].sum(dim=1).cuda() if p_token_lens is None else p_token_lens
+        _eos_penalty[(Logits.shape[0],)](
+            Logits,
+            req_idxs,
+            p_token_lens,
+            exponential_decay_length_penalties,
+            length_penalty_idx,
+            eos_ids,
+            mask_eos_reqs,
+            Logits.stride(0),
+            num_warps=num_warps,
+            EOS_ID_NUM=eos_ids.shape[0],
+        )
+    return
+
+
+if __name__ == "__main__":
+    from .apply_penalty import apply_penalty
+
+    bs = 200
+    vocab_size = 150000
+    p_tokens = 2000
+    repseats = 1000
+    req_idxs = torch.arange(bs).cuda()
+    logits = torch.randn((bs, vocab_size), dtype=torch.float32).cuda()
+    logits2 = logits.clone()
+
+    presence_penalty = torch.randn((bs,), dtype=torch.float32).cuda() + 1e-5
+    freqency_penalty = torch.randn((bs,), dtype=torch.float32).cuda()
+    repetition_penalty = torch.randn((bs,), dtype=torch.float32).cuda()
+    exponential_decay_length_penalties = torch.rand(bs).cuda()
+    eos_ids = torch.tensor([999]).cuda()
+
+    p_seq_len = torch.cat([torch.tensor([0]), torch.randint(1, p_tokens, (bs,))]).cuda()
+    p_token_ids = torch.randint(0, vocab_size, (p_seq_len.sum(),)).cuda()
+    i = 0
+    for s_l in p_seq_len[1:]:
+        p_token_ids[i : i + s_l] = torch.randperm(vocab_size)[:s_l]
+        i += s_l
+    p_token_counts = torch.randint(1, repseats, (p_seq_len.sum(),)).cuda()
+    p_cumsum_seq_len = p_seq_len.cumsum(dim=0).cuda()
+    p_token_vocabs = torch.zeros((bs, vocab_size), dtype=torch.int16).cuda()
+    i = 0
+    b = 0
+    for token_id, count in zip(p_token_ids, p_token_counts):
+        p_token_vocabs[b][token_id] = count
+        i += 1
+        if i == p_seq_len[b + 1]:
+            b += 1
+            i = 0
+
+    p_token_lens = p_token_vocabs.sum(dim=1).cuda()
+    length_penalty_idx = torch.randint(0, p_tokens, (bs,)).cuda()
+    len_idx = torch.tensor([max(p_token_lens[i] - length_penalty_idx[i], 0) for i in range(bs)]).cuda()
+    mask_eos_reqs = torch.randint(1, p_tokens, (bs,)).cuda()
+    mask_bool = torch.tensor([p_token_lens[i] < mask_eos_reqs[i] for i in range(bs)]).cuda()
+
+    fn1 = lambda: apply_penalty(
+        logits,
+        presence_penalty,
+        freqency_penalty,
+        repetition_penalty,
+        p_token_ids,
+        p_token_counts,
+        p_cumsum_seq_len,
+        exponential_decay_length_penalties,
+        len_idx,
+        eos_ids,
+        mask_bool,
+    )
+
+    fn2 = lambda: apply_penalty_cache(
+        logits2,
+        req_idxs,
+        presence_penalty,
+        freqency_penalty,
+        repetition_penalty,
+        p_token_vocabs,
+        p_token_lens,
+        exponential_decay_length_penalties,
+        length_penalty_idx,
+        eos_ids,
+        mask_eos_reqs,
+    )
+    fn1()
+    fn2()
+    cos = F.cosine_similarity(logits, logits2).mean()
+    print("cos =", cos)
+    assert torch.allclose(logits, logits2, atol=1e-2, rtol=0)
+
+    ms1 = triton.testing.do_bench(fn1)
+    ms2 = triton.testing.do_bench(fn2)
+    print("ms1 =", ms1, "ms2 =", ms2)

lightllm/common/req_manager.py

@@ -58,6 +58,7 @@ def __init__(self, max_request_num, max_sequence_length, mem_manager: MemoryMana
             (max_request_num + 1, max_sequence_length), dtype=torch.int32, device="cuda"
         )
         self.mem_manager = mem_manager
+        self.req_sample_parms_manager = None
         self.max_request_num = max_request_num
         self.HOLD_REQUEST_ID = max_request_num
 
@@ -67,6 +68,8 @@ def alloc(self):
     def free(self, free_req_indexes: List[int], free_token_index):
         for req_index in free_req_indexes:
             self.req_list.free(req_index)
+        if self.req_sample_parms_manager is not None:
+            self.req_sample_parms_manager.p_token_vocabs[free_req_indexes] = 0
 
         if self.req_list.is_all_free():
             logger.debug(f"freed all request size {self.req_list.can_alloc_size}")

lightllm/server/router/model_infer/infer_batch.py

@@ -21,6 +21,52 @@
 logger = init_logger(__name__)
 
 
+class ReqSampleParmsManager:
+    def __init__(self, max_request_num, vocab_size):
+        self.presence_penalties = torch.empty(max_request_num, dtype=torch.float, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.frequency_penalties = torch.empty(max_request_num, dtype=torch.float, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.repetition_penalties = torch.empty(max_request_num, dtype=torch.float, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.exponential_decay_length_penalties = torch.empty(
+            max_request_num, dtype=torch.float, device="cpu", pin_memory=True
+        ).cuda(non_blocking=True)
+        self.temperatures = torch.empty(max_request_num, dtype=torch.float, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.top_ps = torch.empty(max_request_num, dtype=torch.float, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.top_ks = torch.empty(max_request_num, dtype=torch.int32, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.length_penalty_idx = torch.empty(max_request_num, dtype=torch.int32, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.mask_eos_reqs = torch.empty(max_request_num, dtype=torch.int32, device="cpu", pin_memory=True).cuda(
+            non_blocking=True
+        )
+        self.p_token_vocabs = torch.zeros(
+            (max_request_num, vocab_size), dtype=torch.int16, device="cpu", pin_memory=True
+        ).cuda(non_blocking=True)
+
+    def set_sample_params(self, req_idx, sampling_param):
+        self.presence_penalties[req_idx] = sampling_param.shm_param.presence_penalty
+        self.frequency_penalties[req_idx] = sampling_param.shm_param.frequency_penalty
+        self.repetition_penalties[req_idx] = sampling_param.shm_param.repetition_penalty
+        tpl = sampling_param.shm_param.exponential_decay_length_penalty.to_tuple()
+        self.exponential_decay_length_penalties[req_idx] = tpl[1]
+        self.temperatures[req_idx] = sampling_param.shm_param.temperature
+        self.top_ps[req_idx] = sampling_param.shm_param.top_p
+        self.top_ks[req_idx] = sampling_param.shm_param.top_k
+        self.length_penalty_idx[req_idx] = tpl[0]
+        self.mask_eos_reqs[req_idx] = sampling_param.shm_param.min_new_tokens - 1
+
+
 @dataclass
 class InferenceContext:
     req_manager: ReqManager = None  # gpu 请求管理
@@ -36,6 +82,8 @@ class InferenceContext:
     def register(
         self, req_manager: ReqManager, radix_cache: RadixCache, shm_req_manager: ShmReqManager, vocab_size: int
     ):
+        if os.getenv("ENABLE_REQ_PARAM_CACHE", False):
+            req_manager.req_sample_parms_manager = ReqSampleParmsManager(req_manager.max_request_num, vocab_size)
         self.req_manager = req_manager
         self.radix_cache = radix_cache
         self.shm_req_manager = shm_req_manager
@@ -55,7 +103,6 @@ def get_overlap_stream(self) -> torch.cuda.Stream:
     def add_reqs(self, requests: List[Tuple[int, int, Any, int]], init_req_obj=True):
         request_ids = []
         for r in requests:
-
             r_id, r_index, multimodal_params, _ = r
             if r_id not in self.requests_mapping.keys():
                 r_obj = InferReq(
@@ -264,6 +311,16 @@ def init_all(self):
             self.shm_req.link_prompt_ids_shm_array()
             self.shm_req.link_logprobs_shm_array()
             self.sampling_param: InferSamplingParams = InferSamplingParams(self.shm_req, self.vocab_size)
+
+            if os.getenv("ENABLE_REQ_PARAM_CACHE", False):
+                g_infer_context.req_manager.req_sample_parms_manager.set_sample_params(
+                    self.req_idx, self.sampling_param
+                )
+                if self.sampling_param.shm_param.input_penalty:
+                    dct = collections.Counter(self.shm_req.get_prompt_ids())
+                    for idx, count in dct.items():
+                        g_infer_context.req_manager.req_sample_parms_manager.p_token_vocabs[self.req_idx][idx] = count
+
             if self.sampling_param.shm_param.input_penalty:
                 self.out_token_id_count = collections.Counter(self.shm_req.get_prompt_ids())
             else:

lightllm/server/router/model_infer/mode_backend/base_backend.py

@@ -226,7 +226,7 @@ def _get_classed_reqs(self, req_ids: List[int], no_decode: bool = False, strict_
                 prefill_reqs.append(req_obj)
                 continue
 
-            is_decode = req_obj.cur_kv_len + 1 == req_obj.get_cur_total_len()
+            is_decode = req_obj.get_output_len() > 0
 
             if not is_decode:
                 prefill_reqs.append(req_obj)