@tduncan / Nvidia-Optimization-Stack-VLLM
Code Commits Issues Pulls Cloud Agent CI/CD
main
main
py 303 lines 10.4 KB
Raw
1 # SPDX-License-Identifier: Apache-2.0
2 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3 from typing import TYPE_CHECKING
4
5 import numpy as np
6 import torch
7
8 from vllm.sampling_params import SamplingParams
9 from vllm.triton_utils import tl, triton
10 from vllm.utils.math_utils import cdiv
11 from vllm.utils.torch_utils import async_tensor_h2d
12 from vllm.v1.worker.gpu.buffer_utils import UvaBackedTensor
13 from vllm.v1.worker.gpu.sample.logits_processor.interface import (
14 LogitsContext,
15 LogitsProcessor,
16 LogitsProcRequestState,
17 )
18
19 if TYPE_CHECKING:
20 from vllm.config import VllmConfig
21
22
23 class PenaltiesState(LogitsProcessor):
24 def __init__(self, vllm_config: "VllmConfig", req_states: LogitsProcRequestState):
25 self.req_states = req_states
26 self.device = req_states.device
27 max_num_reqs = req_states.max_num_reqs
28 vocab_size = req_states.vocab_size
29
30 self.repetition_penalty = UvaBackedTensor(max_num_reqs, dtype=torch.float32)
31 self.frequency_penalty = UvaBackedTensor(max_num_reqs, dtype=torch.float32)
32 self.presence_penalty = UvaBackedTensor(max_num_reqs, dtype=torch.float32)
33 self.use_penalty = np.zeros(max_num_reqs, dtype=bool)
34
35 # Initialize repetition penalty manually because 0 is an invalid value for it.
36 self.repetition_penalty.np.fill(1.0)
37 self.repetition_penalty.copy_to_uva()
38
39 # Statistics for penalties.
40 self.prompt_bin_mask = torch.zeros(
41 max_num_reqs, cdiv(vocab_size, 32), dtype=torch.int32, device=self.device
42 )
43 # TODO(woosuk): This tensor is rarely used but can be very large, taking up
44 # GBs of GPU memory. Optimize the memory usage.
45 self.output_bin_counts = torch.zeros(
46 max_num_reqs, vocab_size, dtype=torch.int32, device=self.device
47 )
48
49 self._new_penalties_reqs: list[int] = []
50
51 def add_request(self, req_idx: int, sampling_params: SamplingParams) -> bool:
52 self.repetition_penalty.np[req_idx] = sampling_params.repetition_penalty
53 self.frequency_penalty.np[req_idx] = sampling_params.frequency_penalty
54 self.presence_penalty.np[req_idx] = sampling_params.presence_penalty
55
56 do_penalty = use_penalty(sampling_params)
57 self.use_penalty[req_idx] = do_penalty
58 if do_penalty:
59 self._new_penalties_reqs.append(req_idx)
60 return do_penalty
61
62 def apply_staged_writes(self) -> None:
63 if self._new_penalties_reqs:
64 idx_mapping = async_tensor_h2d(
65 self._new_penalties_reqs, dtype=torch.int32, device=self.device
66 )
67
68 prefill_lens = self.req_states.prefill_len.np[self._new_penalties_reqs]
69 max_prefill_len = int(prefill_lens.max())
70 bincount(
71 idx_mapping,
72 self.req_states.all_token_ids.gpu,
73 self.req_states.prompt_len.gpu,
74 self.req_states.prefill_len.gpu,
75 self.prompt_bin_mask,
76 self.output_bin_counts,
77 max_prefill_len,
78 )
79 self._new_penalties_reqs.clear()
80
81 self.repetition_penalty.copy_to_uva()
82 self.frequency_penalty.copy_to_uva()
83 self.presence_penalty.copy_to_uva()
84
85 def apply(self, logits: torch.Tensor, ctx: LogitsContext) -> torch.Tensor:
86 if not np.any(self.use_penalty[ctx.idx_mapping_np]):
87 # No request uses penalties. Skip the kernel launch.
88 return logits
89
90 apply_penalties(
91 logits,
92 ctx.expanded_idx_mapping,
93 ctx.input_ids,
94 ctx.expanded_local_pos,
95 self.repetition_penalty.gpu,
96 self.frequency_penalty.gpu,
97 self.presence_penalty.gpu,
98 self.prompt_bin_mask,
99 self.output_bin_counts,
100 )
101 return logits
102
103
104 @triton.jit
105 def _penalties_kernel(
106 logits_ptr,
107 logits_stride,
108 expanded_idx_mapping_ptr,
109 token_ids_ptr,
110 expanded_local_pos_ptr,
111 repetition_penalty_ptr,
112 frequency_penalty_ptr,
113 presence_penalty_ptr,
114 prompt_bin_mask_ptr,
115 prompt_bin_mask_stride,
116 output_bin_counts_ptr,
117 output_bin_counts_stride,
118 vocab_size,
119 BLOCK_SIZE: tl.constexpr,
120 ):
121 token_idx = tl.program_id(0).to(tl.int64)
122 req_state_idx = tl.load(expanded_idx_mapping_ptr + token_idx)
123 rep_penalty = tl.load(repetition_penalty_ptr + req_state_idx)
124 freq_penalty = tl.load(frequency_penalty_ptr + req_state_idx)
125 pres_penalty = tl.load(presence_penalty_ptr + req_state_idx)
126
127 use_rep_penalty = rep_penalty != 1.0
128 use_freq_penalty = freq_penalty != 0.0
129 use_pres_penalty = pres_penalty != 0.0
130 use_penalty = use_rep_penalty or use_freq_penalty or use_pres_penalty
131 if not use_penalty:
132 # Early return to avoid loading logits.
133 return
134
135 block_idx = tl.program_id(1)
136 block = block_idx * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
137 mask = block < vocab_size
138 logits = tl.load(logits_ptr + token_idx * logits_stride + block, mask=mask)
139 logits = logits.to(tl.float32)
140
141 base_output_counts = tl.load(
142 output_bin_counts_ptr + req_state_idx * output_bin_counts_stride + block,
143 mask=mask,
144 other=0,
145 )
146
147 # Accumulate draft token counts from previous positions directly into
148 # output_bin_counts (preserves its native tensor layout, avoiding an
149 # expensive shared-memory layout conversion after the loop).
150 pos = tl.load(expanded_local_pos_ptr + token_idx)
151 start_idx = token_idx - pos
152 output_bin_counts = base_output_counts
153 for prev_pos in tl.range(pos):
154 prev_token = tl.load(token_ids_ptr + start_idx + prev_pos + 1)
155 token_match = block == prev_token
156 output_bin_counts = output_bin_counts + token_match.to(tl.int32)
157 output_bin_mask = output_bin_counts > 0
158
159 # Apply repetition penalties.
160 if use_rep_penalty:
161 packed_block = block_idx * BLOCK_SIZE // 32 + tl.arange(0, BLOCK_SIZE // 32)
162 packed_mask = tl.load(
163 prompt_bin_mask_ptr + req_state_idx * prompt_bin_mask_stride + packed_block,
164 mask=packed_block < tl.cdiv(vocab_size, 32),
165 other=0,
166 )
167 prompt_bin_mask = (packed_mask[:, None] >> (tl.arange(0, 32)[None, :])) & 1
168 prompt_bin_mask = prompt_bin_mask.to(tl.int1)
169 prompt_bin_mask = prompt_bin_mask.reshape(BLOCK_SIZE)
170
171 # If token appears in prompt or output, apply, otherwise use 1.0 for no-op.
172 scale = tl.where(prompt_bin_mask | output_bin_mask, rep_penalty, 1.0)
173 # If logits are positive, divide by penalty, otherwise multiply by penalty.
174 logits *= tl.where(logits > 0, 1.0 / scale, scale)
175
176 # Apply frequency penalties.
177 logits -= freq_penalty * output_bin_counts
178 # Apply presence penalties.
179 logits -= pres_penalty * output_bin_mask
180 # Store back to logits.
181 tl.store(logits_ptr + token_idx * logits_stride + block, logits, mask=mask)
182
183
184 def apply_penalties(
185 logits: torch.Tensor,
186 expanded_idx_mapping: torch.Tensor,
187 token_ids: torch.Tensor,
188 expanded_local_pos: torch.Tensor,
189 repetition_penalty: torch.Tensor,
190 frequency_penalty: torch.Tensor,
191 presence_penalty: torch.Tensor,
192 prompt_bin_mask: torch.Tensor,
193 output_bin_counts: torch.Tensor,
194 ) -> None:
195 num_tokens, vocab_size = logits.shape
196 BLOCK_SIZE = 8192
197 num_blocks = triton.cdiv(vocab_size, BLOCK_SIZE)
198 _penalties_kernel[(num_tokens, num_blocks)](
199 logits,
200 logits.stride(0),
201 expanded_idx_mapping,
202 token_ids,
203 expanded_local_pos,
204 repetition_penalty,
205 frequency_penalty,
206 presence_penalty,
207 prompt_bin_mask,
208 prompt_bin_mask.stride(0),
209 output_bin_counts,
210 output_bin_counts.stride(0),
211 vocab_size,
212 BLOCK_SIZE=BLOCK_SIZE,
213 )
214
215
216 @triton.jit
217 def _bincount_kernel(
218 expanded_idx_mapping_ptr,
219 all_token_ids_ptr,
220 all_token_ids_stride,
221 prompt_len_ptr,
222 prefill_len_ptr,
223 prompt_bin_mask_ptr,
224 prompt_bin_mask_stride,
225 output_bin_counts_ptr,
226 output_bin_counts_stride,
227 BLOCK_SIZE: tl.constexpr,
228 ):
229 token_idx = tl.program_id(0)
230 block_idx = tl.program_id(1)
231 req_state_idx = tl.load(expanded_idx_mapping_ptr + token_idx)
232
233 prefill_len = tl.load(prefill_len_ptr + req_state_idx)
234 if block_idx * BLOCK_SIZE >= prefill_len:
235 return
236
237 prompt_len = tl.load(prompt_len_ptr + req_state_idx)
238 block = block_idx * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
239 if block_idx * BLOCK_SIZE < prompt_len:
240 mask = block < prompt_len
241 prompt_tokens = tl.load(
242 all_token_ids_ptr + req_state_idx * all_token_ids_stride + block, mask=mask
243 )
244 idx = prompt_tokens // 32
245 bit_idx = prompt_tokens % 32
246 bit = tl.full((BLOCK_SIZE,), 1, tl.int32) << bit_idx
247 tl.atomic_or(
248 prompt_bin_mask_ptr + req_state_idx * prompt_bin_mask_stride + idx,
249 bit,
250 mask=mask,
251 )
252
253 if (block_idx + 1) * BLOCK_SIZE >= prompt_len:
254 mask = block < prefill_len
255 mask &= block >= prompt_len
256 output_tokens = tl.load(
257 all_token_ids_ptr + req_state_idx * all_token_ids_stride + block, mask=mask
258 )
259 tl.atomic_add(
260 output_bin_counts_ptr
261 + req_state_idx * output_bin_counts_stride
262 + output_tokens,
263 1,
264 mask=mask,
265 )
266
267
268 def bincount(
269 expanded_idx_mapping: torch.Tensor,
270 all_token_ids: torch.Tensor,
271 prompt_len: torch.Tensor,
272 prefill_len: torch.Tensor,
273 prompt_bin_mask: torch.Tensor,
274 output_bin_counts: torch.Tensor,
275 max_prefill_len: int,
276 ) -> None:
277 # Use index_fill_ instead of `tensor[idx] = 0` to avoid sync.
278 idx_long = expanded_idx_mapping.long()
279 prompt_bin_mask.index_fill_(0, idx_long, 0)
280 output_bin_counts.index_fill_(0, idx_long, 0)
281 num_tokens = expanded_idx_mapping.shape[0]
282 BLOCK_SIZE = 1024
283 num_blocks = triton.cdiv(max_prefill_len, BLOCK_SIZE)
284 _bincount_kernel[(num_tokens, num_blocks)](
285 expanded_idx_mapping,
286 all_token_ids,
287 all_token_ids.stride(0),
288 prompt_len,
289 prefill_len,
290 prompt_bin_mask,
291 prompt_bin_mask.stride(0),
292 output_bin_counts,
293 output_bin_counts.stride(0),
294 BLOCK_SIZE=BLOCK_SIZE,
295 )
296
297
298 def use_penalty(sampling_params: SamplingParams) -> bool:
299 return (
300 sampling_params.repetition_penalty != 1.0
301 or sampling_params.frequency_penalty != 0.0
302 or sampling_params.presence_penalty != 0.0
303 )