add eagle

Signed-off-by: Zhi Yiliu <2584074296@qq.com>

Author: tom-jerr

.gitignore

@@ -6,4 +6,5 @@ TaskNotes/
 .ai_cache
 docs/assets
 _ai_summary*.json
-.python
+.python
+.venv

docs/notes/cuda/cuda basic.md

@@ -32,7 +32,39 @@ cudaMemcpyAsync(dst, src, size, cudaMemcpyHostToDevice, stream);
   - 如果放在不同 stream，并且硬件支持 copy engine，那么数据拷贝和计算可以 overlap：
     - stream0: H2D 拷贝下一批
     - stream1: 执行当前批 kernel
+# GPU 异步拷贝
+## Pinned Memory
+Pinned memory 就是 页锁定内存（page-locked memory），是主机内存里一块不会被操作系统随意换页的区域。
 
+如果 host 端是普通 pageable memory：
+
+- CUDA 很可能先偷偷拷到一块内部 staging pinned buffer
+- 再 DMA 这样会多一次 copy
+- 甚至可能退化成近似同步行为
+
+而 pinned memory 的好处是：
+
+- GPU/网卡可以直接 DMA 访问
+- H2D / D2H 拷贝更快
+- 更重要的是 支持真正的异步拷贝
+
+## 异步拷贝
+发起一次数据搬运后，CPU 不必堵塞等待，GPU/拷贝引擎可以在后台执行，和别的计算重叠。
+
+比如 CUDA 里的：
+```cpp
+cudaMemcpyAsync(dst, src, bytes, cudaMemcpyDeviceToHost, stream);
+```
+它表示：
+
+- 把这个 copy 操作插入某个 stream，调用返回时，拷贝可能还没完成
+- 后续可以继续发别的 kernel / copy
+- 最后通过 event / stream synchronize 判断何时完成
+
+## 注意事项
+- 实际使用时会申请一大块内存作为 pinned memory pool 来避免频繁申请释放的开销
+  - 避免频繁 cudaHostAlloc/cudaFreeHost 带来的高开销和抖动
+- 使用时先从 pool 取一个 pinned buffer，再通过 cudaMemcpyAsync 发起异步 H2D/D2H 拷贝，并用 cudaEventRecord 标记完成时刻。CPU 侧等 event 完成后再消费数据，最后把 buffer 放回池中。
 ## Kernel Launch
 - 一个 kernel launch 到 GPU 上发生了什么
   1. CPU 发起 kernel launch。

docs/notes/cuda/torch compile.md

@@ -0,0 +1,138 @@
+---
+
+title: Model Forward 瓶颈以及优化方法
+created: 2026-04-05
+updated:
+tags:
+
+- LLMInference
+description: Model Forward 瓶颈以及优化方法
+# cover: /img/transformer.png
+---
+# Model Forward 瓶颈分析以及优化方法
+这里以 MHA 为例，输入 X 大小为 [B, S, D]，其中 B 是 batch size，S 是序列长度，D 是隐藏层维度。
+- 每一层是一个 Transformer Block，包含一个 MHA 和一个 FFN，中间还有 Norm 和残差连接。
+
+MHA 的计算过程如下：
+1. QKV: Wq, Wk, Wv 分别是大小为 [D, D] 的权重矩阵，计算 Q = X * Wq, K = X * Wk, V = X * Wv，得到 Q, K, V 大小为 [B, S, D] -> [B, S, H, D_h] -> [B, H, S, D_h]
+2. Attention Scores: 计算 Attention Scores = Q * K^T / sqrt(D)，得到大小为 [B, H, S, S]。
+3. softmax: 对 Attention Scores 进行 softmax，得到 Attention Weights 大小为 [B, H, S, S]。
+4. score * V: 计算 Output = Attention Weights * V，得到大小为 [B, H, S, D_h] -> [B, S, D]。
+5. 输出 Output 大小为 [B, S, D]。
+
+FFN 的计算过程如下：
+1. Linear1: W1 大小为 [D, D_ffn]，计算 Output1 = X * W1，得到大小为 [B, S, D_ffn]。
+2. Activation: 对 Output1 进行激活函数（如 ReLU），得到大小为 [B, S, D_ffn]。
+3. Linear2: W2 大小为 [D_ffn, D]，计算 Output2 = Output1 * W2，得到大小为 [B, S, D]。
+4. Linear3: W3 大小为 [D, D]，计算 Output3 = X * W3，得到大小为 [B, S, D]。
+5. 输出 Output 大小为 [B, S, D]。
+
+## Arithmetic Intensity 分析(FP16)
+### Prefill 阶段
+#### MHA
+QKV 计算每个都是一个 GEMM 的算术强度为:
+$$
+AI = \frac{2 * B * S * D * D}{2*(B * S * D(input) + D * D(weight) + B * S * D(output))} = \frac{2 * B * S * D^2}{2*(2 * B * S * D + D^2)} \approx \frac{D}{2 + \frac{D}{BS}}
+$$
+
+Attention Scores 计算的算术强度为:
+- QK^T 的算术强度为:
+$$
+AI = \frac{2 * B * H * S^2 * D_h}{B * S * D + B * S * D + B * H * S^2} = \frac{2 * B * H * S^2 * D_h}{2*(2 * B * S * D + B * H * S^2)} \approx \frac{SD_h}{2D_h + S}
+$$
+- softmax 的算术强度为: 基本上都是 Memory-Bound，算术强度非常低。
+$$
+AI = \frac{B * H * S*(5S - 2)}{2*(B * H * S^2 + B * H * S^2)} \approx \frac{5}{4}
+$$
+
+- score * V 的算术强度为:
+$$
+AI = \frac{2 * B * H * S^2 * D_h}{2*(B * H * S^2 + B * S * D + B * S * D)} \approx \frac{D_h}{1 + \frac{2D_h}{S}}
+$$
+
+
+#### FFN
+- Linear1 的算术强度为: 一般 D_ffn 是 D 的 4 倍。
+$$
+AI = \frac{2 * B * S * D * D_{ffn}}{2*(B * S * D + D * D_{ffn} + B * S * D_{ffn})} \approx \frac{D_{ffn}}{5 + \frac{D_{ffn}}{BS}}
+$$
+- Activation 的算术强度为: 基本上都是 Memory-Bound，算术强度非常低。
+- Linear2 的算术强度为:
+$$
+AI = \frac{2 * B * S * D_{ffn} * D}{2*(B * S * D_{ffn} + D_{ffn} * D + B * S * D)} \approx \frac{D}{1 + \frac{D}{D_{ffn}} + \frac{D}{BS}}
+$$
+- Linear3 的算术强度为:
+$$
+AI = \frac{2 * B * S * D^2}{2*(B * S * D + D * D + B * S * D)} \approx \frac{D}{2 + \frac{D}{BS}}
+$$
+
+### Decode 阶段
+- 此时每一个 Q 长度都为 1，K 和 V 的长度是 S（历史键值对 + 当前输入）
+#### MHA
+
+Attention Scores 计算的算术强度为:
+- QK^T 的算术强度为:
+$$
+AI = \frac{2 * B * H * S * 1 * D_h}{2 *(B * 1 * D + B * S * D + B * H * S^2)}  \approx \frac{D_h}{D_h + S + \frac{D_h}{S}}
+$$
+- softmax 的算术强度为: 基本上都是 Memory-Bound，算术强度非常低。
+$$
+AI = \frac{B * H * S*(5S - 2)}{2*(B * H * S^2 + B * H * S^2)} \approx \frac{5}{4}
+$$
+
+- score * V 的算术强度为:
+$$
+AI = \frac{2 * B * H * S * 1 * D_h}{2*(B * H * S * 1 + B * S * D + B * 1 * D)} \approx \frac{D_h}{1 + D_h + \frac{D_h}{S}}
+$$
+
+### 分析
+- Prefill 阶段 S 比较大，算术强度相对较高，主瓶颈常在 GEMM 算力利用率 和 attention 中间张量 IO
+  - QK^T 产生 [B, H, S, S] 的中间张量，softmax 又要读写这个张量，score·V 再读一次
+- Decode 阶段，每步只有一个 query token，Q 长度是 1，K/V 长度是历史长度 S，每生成一个 token 都要读整段历史 KV cache，瓶颈一般在KV cache 读取带宽以及 KV cache 的布局 / 碎片 / 页访问局部性
+  - FFN / projection 虽然 FLOPs 大，但因 S=1 shape 很小，常常也吃不满算力
+
+## 优化方法
+### Batching
+- 对 Prefill 阶段提升不是很大，因为 S 已经比较大了已经是 Compute-Bound 了，提升算力利用率的空间有限。
+- 对 Decode 阶段提升很大，可以显著提升算力利用率，减少每个 token 的平均 latency。
+
+### GQA
+- 每个 KV head 为多个 query head 服务，减少了需要维护和加载的 KV cache 的数量，降低了内存带宽压力。
+  - 因为共享 K head 后，一个 K tile 可以服务多个 Q heads：那么实际 HBM 读取会下降，cache locality 更好。这会提高实际有效 AI
+- Decode AI，输入 [B, 1, D]，输出 [B,1, D]，算术强度为:
+  - 实际上与 H_q 和 H_k 的比值相关，
+$$
+AI = \frac{2 * B * H_q * S * 1 * D_h}{2*(B * 1 * H_q * D_h + B * S * H_k * D_h + B * S * H_k * D_h + B * 1 * H_q * D_h)} \approx \frac{H_q}{2(H_k + \frac{H_q}{S})}
+$$
+  - GQA 的理论 AI 主要用来判断 decode attention 是否仍然受 KV 带宽限制，以及减小 H_k 还有没有收益。长上下文下这个 AI 近似与 H_q/H_k 成正比
+  - GQA 的本质价值是提升每次 K/V 读取的复用度。
+- 调优：如果没有达到预期，通常从几个方向排查：
+  1. 是否有额外 gather/scatter、transpose 或中间张量导致实际 bytes 远大于理论；
+  2. 共享 K/V 是否真的在片上复用了，还是不同 Q heads 仍然重复从 HBM 加载；
+  3. 为了复用而增加的寄存器和 shared memory 占用是否把 occupancy 打低了；
+  4. paged KV 的碎片和非连续访问是否吞掉了 GQA 带来的带宽收益；
+  5. 最后还要确认端到端瓶颈有没有转移到 sampling、通信或 CPU-GPU 同步。
+
+### Speculative Decoding
+- 可以理解为同样是为了提高 decode 阶段的算力利用率，但它是通过增加并行度（同时生成多个 token，即增大 S）来提升理论 AI，同时减少 KV Cache 的重复访问提高有效 AI
+
+
+### Flash Attention
+- 并没有提升理论 AI，但通过减少中间张量的读写来提升实际 AI。
+- 通过 fuse Attention 的多个步骤（QK^T、softmax、score·V）来减少中间张量的读写，降低内存带宽需求，提高实际 AI。
+  - 每次加载一个 Q/K/V tile 后，直接计算对应的 Attention Scores 和 score·V 的部分结果，避免了生成完整的 [B, H, S, S] Attention Scores 张量。
+  ![](img/flash_attention_algo.png)
+
+
+
+
+
+
+
+
+
+
+
+
+
+

docs/notes/llm_inference/asynccopy.md

@@ -1,3 +1,4 @@
+
 # CUDA Graph
 ## 原理
 CUDA Graph 将一段 GPU kernel 序列录制为静态 DAG，之后只需一次 CPU launch 即可重放整个计算流，以此消除逐个 kernel launch 的 CPU 开销。在此基础上，我们更进一步理解 CUDA Graph 的一些核心机制。
@@ -72,4 +73,50 @@ The full FX graph is passed to Inductor; Inductor's own scheduler performs parti
 ### 流程
 - 对每个非 splitting 子图，创建 PiecewiseBackend 并编译所有 ranges（包括 symbolic range 和 static sizes）
 - PiecewiseBackend 在初始化时调用 compile_all_ranges()，对每个 compile_range 提前编译好对应的 runnable。运行时根据 runtime_shape 查找对应 range 的 runnable 并执行
-- 非 splitting 子图的 PiecewiseBackend 最终被 wrap_with_cudagraph_if_needed() 包裹成 CUDA Graph wrapper
+- 非 splitting 子图的 PiecewiseBackend 最终被 wrap_with_cudagraph_if_needed() 包裹成 CUDA Graph wrapper
+# torch.compile
+```shell
+Python eager 代码
+→ TorchDynamo 截获 Python frame / 字节码
+→ 抽取 FX Graph + 记录 guards(这份编译结果成立的前提条件)
+→ AOTAutograd（训练场景）做前反向图分解
+→ TorchInductor 做图级优化、fusion、调度、代码生成
+→ 后端落到 Triton / CPP / ATen 等
+→ Triton 再把 kernel 编译成设备可执行代码
+→ 运行时按 guards 检查是否可复用旧编译结果
+```
+- Dynamo 更像一个 Python 层图抽取器 + 特化入口，不是最后生成 GPU kernel 的地方。官方文档也明确说：每个 frame 都会尝试编译，并把**编译结果缓存在 code object 上；如果后续调用不满足之前的 guards，就会发生 guard failure，然后重新编译**
+
+## compile cache
+### 第一层：Dynamo 的 frame / guard 级缓存
+
+缓存的是：
+- 某段 Python frame 对应的已编译结果
+- 它配套的 guards
+
+只要这次调用还满足旧 guards，就直接复用，不重新抓图/编译。官方文档明确说它会把编译结果缓存在 code object 上，并在 guard failure 时重新编译。
+
+### 第二层：Inductor 的图级缓存
+
+缓存的是：
+
+- FX Graph / 规范化后的图表示
+- 对应生成过的后端代码
+- 一些中间 IR 或编译结果索引
+
+PyTorch 官方 compile caching recipe 里明确提到，默认有本地磁盘缓存，里面包括 FXGraphCache 等模块化缓存。
+
+### 第三层：Triton kernel 编译缓存
+
+缓存的是：
+
+- Triton kernel 源/IR 对应的设备代码
+- 和 launch 相关的元数据
+- autotune 结果或编译 key 相关信息
+
+这一层更接近“传统 JIT kernel cache”。
+## 原理
+`torch.compile` 是 PyTorch 提供的一个用于加速模型推理和训练的工具。它通过将 PyTorch 代码转换为更高效的中间表示（IR），并利用各种优化技术来提升性能。其核心原理包括以下几个方面：
+- Dynamo: 拦截 Python 执行代码，提取 tensor op 构成静态 IR 图，跳过 Python 解释器
+- Inductor: 算子融合 + kernel 合并 + 内存优化
+- CUDAGraph: 减少 CPU 调度开销，一次录制多次执行

docs/notes/llm_inference/img/flash_attention_algo.png

@@ -1,3 +1,15 @@
+---
+
+title: Speculative Decoding
+created: 2026-04-5
+update:
+comments: true
+katex: true
+tags:
+
+- LLMInference
+
+---
 # Speculative Decoding
 
 
@@ -140,8 +152,11 @@ $$
 1. verify 完成后，SGLang 会真正把接受的 token 写入 req.output_ids，更新 grammar、kv_committed_len、seq_lens
      - 释放未接受分支占掉的 KV 槽；topk>1 时还会做 cache compaction，把保留下来的 accepted path 挪到连续位置。
 2. verify 得到的 prefict 构造 forward batch，draft 模型沿着这次真正接受的 token 序列做一次 extend，重新得到下一轮要用的 topk_p/topk_index/hidden_states。
+
 3. 构造下一次的 EagleDraftInput，主要是把这轮 verify 接受的 token(最后一个被验证成功的 token) 以及得到的 topk_p/topk_index/hidden_states。这一步非常关键，它让 draft 始终跟 target 已确认的前缀同步。
 
+> [!IMPORTANT]
+> 这里是因为 EAGLE 本身的性质决定的，它需要 target model 的 feature 和 sample 后的 token 作为 draft model 的输入来预测 draft token，所以每次 draft->verify->sampling 之后，我们都需要再次推进 draft，即用 target model 刚刚得到的所有 token 的 hidden_state 以及采样得到的 bonus token 来恢复 KV Cache 以供下一轮 draft 使用。
 ## 工程细节 SGLang
 ### Mamba Cache Problem
 - 普通 Transformer 在 speculative verify 后，只需要把接受的 token 写回 target KV cache 就够了。
@@ -158,8 +173,8 @@ $$
     - 一次把 mamba_steps_to_track scatter 到 mamba_track_indices，让 prefix-cache 的 snapshot 落在正确的 tracking 点上。
 
 **Important** 
-- Mamba cache 不是单纯 token->KV 的映射，而是“序列状态快照”。
-- speculative verify 会先算出一串候选，再只接受其中一部分；真正该缓存的 Mamba state，不一定是 verify 最后一步，而是“accepted 路径里跨过 track interval 的那一步”。
+- Mamba cache 不是单纯 token->KV 的映射，而是序列状态快照。
+- speculative verify 会先算出一串候选，再只接受其中一部分；真正该缓存的 Mamba state，不一定是 verify 最后一步，而是 accepted 路径里跨过 track interval 的那一步。
 - verify 路径和 extend 路径用的 tracking 语义不一样，所以这里只保留 mamba_track_indices，并显式把 mamba_track_mask、mamba_track_seqlens 置空；否则会把 extend 阶段遗留的 metadata 误带进 verify