@tduncan / Nvidia-Optimization-Stack-VLLM Mirror · read-only / commits / 899541bdb1
Fork 0

A high-throughput and memory-efficient inference and serving engine for LLMs

Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git

Synced Nov 10, 07:45 UTC

[doc] fix broken links (#32158)

Signed-off-by: minimAluminiumalism <caixuesen@outlook.com>

XlKsyt committed Jan 12, 2026 at 18:18 UTC 899541bdb1e42762c17a8cb5cf5bf2800ef63811
1 file changed +7 -21
docs/design/paged_attention.md
+7 -21
@@ -139,18 +139,14 @@ token data.
139 const scalar_t* q_ptr = q + seq_idx * q_stride + head_idx * HEAD_SIZE;
140 ```
141
142 -<p align="center">
143 - <img src="../assets/design/paged_attention/query.png" alt="query" width="70%" />
144 -</p>
142 +![query](../assets/design/paged_attention/query.png)
143
144 Each thread defines its own `q_ptr` which points to the assigned
145 query token data on global memory. For example, if `VEC_SIZE` is 4
146 and `HEAD_SIZE` is 128, the `q_ptr` points to data that contains
147 total of 128 elements divided into 128 / 4 = 32 vecs.
148
151 -<p align="center">
152 - <img src="../assets/design/paged_attention/q_vecs.png" alt="q_vecs" width="70%" />
153 -</p>
149 +![q_vecs](../assets/design/paged_attention/q_vecs.png)
150
151 ```cpp
152 __shared__ Q_vec q_vecs[THREAD_GROUP_SIZE][NUM_VECS_PER_THREAD];
@@ -187,9 +183,7 @@ key token at different iterations. As shown above, that `k_ptr`
183 points to key token data based on `k_cache` at assigned block,
184 assigned head and assigned token.
185
190 -<p align="center">
191 - <img src="../assets/design/paged_attention/key.png" alt="key" width="70%" />
192 -</p>
186 +![key](../assets/design/paged_attention/key.png)
187
188 The diagram above illustrates the memory layout for key data. It
189 assumes that the `BLOCK_SIZE` is 16, `HEAD_SIZE` is 128, `x` is
@@ -202,9 +196,7 @@ iterations. Inside each rectangle, there are a total 32 vecs (128
196 elements for one token) that will be processed by 2 threads (one
197 thread group) separately.
198
205 -<p align="center">
206 - <img src="../assets/design/paged_attention/k_vecs.png" alt="k_vecs" width="70%" />
207 -</p>
199 +![k_vecs](../assets/design/paged_attention/k_vecs.png)
200
201 ```cpp
202 K_vec k_vecs[NUM_VECS_PER_THREAD]
@@ -361,17 +353,11 @@ later steps. Now, it should store the normalized softmax result of
353
354 ## Value
355
364 -<p align="center">
365 - <img src="../assets/design/paged_attention/value.png" alt="value" width="70%" />
366 -</p>
356 +![value](../assets/design/paged_attention/value.png)
357
368 -<p align="center">
369 - <img src="../assets/design/paged_attention/logits_vec.png" alt="logits_vec" width="50%" />
370 -</p>
358 +![logits_vec](../assets/design/paged_attention/logits_vec.png)
359
372 -<p align="center">
373 - <img src="../assets/design/paged_attention/v_vec.png" alt="v_vec" width="70%" />
374 -</p>
360 +![v_vec](../assets/design/paged_attention/v_vec.png)
361
362 Now we need to retrieve the value data and perform dot multiplication
363 with `logits`. Unlike query and key, there is no thread group