Yuchen You

recvReduceSend() reduce-scatter 本地实现 1234 __device__ __forceinline__ void recvReduceSend(intptr_t inpIx, int eltN, bool postOp=false) { genericOp<0, 0, 1, 1, Input, -1>(inpIx, -1, eltN, postOp); // DirectRecv1, DirectSend1, Recv, Send, SrcBuf, DstBuf} direct 这里指的是可以直接访问到 peer 的 gpu 显存，但是这里并不一定是远程 gpu 可能是同一个 server 内部的不同 gpu 设备

整体架构 Kimi Linear 将 KDA 和 MLA 放在 Transformer 的 token-mixing 位置, 并采用 3 个 KDA layers + 1 个 MLA layer 的排列; KDA 用固定大小的 recurrent state 处理历史信息, MLA 则保留 full attention 与 latent KV cache; 每个 token-mixing layer 后仍然连接一个 MoE layer, 整体数据流保持 Norm -> KDA 或 MLA -> residual -> Norm -> MoE -> residual; KDA layer 的输入, 状态更新与输出 生成第 t+1t+1t+1 个 token 时, KDA layer lll 接收两个输入: 当前 token 从上一层传来的 hidden vector ht(l)h_t^{(l)}ht(l)​, 以及本层处理完上一个 token 后留下的 recurrent state St−1(l)S_{t-1}^{(l)}St−1(l)​; 前者沿模型深度 ...

背景 MoE 模型的总参数量虽然很大, 但每个 token 实际只会激活其中少量 experts, 因此 active parameters 远小于总参数量, 这个现象称为 sparse activation; 它让消费级 GPU 具备了计算大规模 MoE 模型的可能性, 因为 GPU 只需要执行当前激活的部分, 不需要同时计算完整的 expert pool; 但是, 完整 expert pool 仍然可能远大于 GPU 显存, 因此需要主要保存在 host memory 中; FreeToken 希望把个人电脑的 GPU, CPU, memory 和 interconnect 组合成一个完整的 serving system; 当需要的 expert 没有命中 GPU cache 时, 系统可以将 expert weights 搬到 GPU 上计算, 也可以将较小的 activation tensor 传到 CPU 侧直接计算; FreeToken 希望感知 CPU, GPU 和 memory 之间的实际 bandwidth, 并据此动态选择执行方式, 从而尽可能利用个人电脑已有的异构 ...

模型生成到一半被发现越狱了, 或者被发现在编事实。这一刻你能做什么? 这份笔记把"检测之后的动作"整理成八类, 梳理 steering 和 rollback 两条主线的谱系、当前 SOTA 和还没被解决的问题。文献覆盖 2023.06 – 2026.08, 约 60 篇。 分类学是我为回答这个问题归纳的, 不是文献里的既有共识。要写进 related work 的话, 建议把 A1/A2 合并成 “continuous intervention”, A4/A5/A6 合并成 “generative rollback”, 更接近现有说法。 1. 先把术语对齐 这两个词在文献里各有好几个名字, 而且它们不是并列关系, 是同一个闭环里的两种不同动作。 Steering 的正式说法主要有三个: activation steering / activation addition: 往残差流里加一个方向向量 representation engineering (RepE): 把这件事当成一个研究纲领 inference-time intervention (ITI): ...

Multi-Head Latent Attention Low-Rank Key-Value Joint Compression DeepseekMoE

📝 到达过程黑话 arrival process (到达过程): 事件在时间轴上一个个冒出来的随机过程 N(t)N(t)N(t): 到时刻 ttt 为止发生了多少次事件 (计数视角) TkT_kTk​: 第 kkk 次事件发生的时刻 (等待视角) λ\lambdaλ: rate, 单位时间平均发生几次 survival function S(t)=P(T>t)S(t) = P(T > t)S(t)=P(T>t): “撑到 ttt 还没出事” 的概率 hazard rate h(t)h(t)h(t): 风险率, 已经活到 ttt 了, 下一瞬间出事的密度 shape parameter kkk: Weibull 里控制 hazard 走势的那个形状参数 123456789Binomial (n 次独立试验, 每次概率 p) ↓ n → ∞, p → 0, np = λt 固定Poisson(λt) "固定一段时间 t, 数一数发生了几次" ↓ 反过来问: 等第 1 次要多久Exponenti ...

Speculative Decoding 的朴素思想 想要加速一个 auto-regressive 模型的推理, 可以先猜测合理的上下文, 再让大模型一次验证多个位置; 比如已经有 prefix, 小模型先猜出 I love cats, 大模型验证后保留 I love, 在 cats 这里拒绝并补出 dogs, 后面基于 cats 生成的内容就不能继续保留; 这里有两个角色: Draft Model 负责快速生成候选, Target Model 负责验证和确定最终输出; 上面的例子是在说明接受前缀和截断的位置, 随机 sampling 下的接受规则还需要下面的概率推导; 本文主要根据四页手稿整理, 保留从 sampling 理论到 SpecInfer, Medusa, EAGLE-1 的顺序; 图片选用原论文架构图, 对手稿中公式或表述的必要修正单独说明; 理论部分对应 Google 的 Fast Inference from Transformers via Speculative Decoding 和 DeepMind 的 Accelerating Large Language M ...

背景 Tokenwise locality Layerwise sensitivity to loss Encoding 和 Decoding Delta Encoding 我们在构建搜索引擎的底层存储编码格式的时候用过一个 Delta Encoding/Decoding 的思路，将对应 token 坐标通过相对位置来存储，可以让大数字变成较小的相对位移来存储来减少存储量；这里用到了非常类似的思路: 对于一个 token 的值 (KV cache)，我们用相对于前面值的变化量来存储，当然在这里需要一个 microbench 来说明这个存储在 token kv 背景下确实可以减少存储量: 如上图所示的测量中可以看出来，这是累积分布，所以 Delta 编码的斜率在靠近 0 点的时候都很大，说明 Delta 存储下可以有效利用小数据尺寸来存储大量的数据以节省空间; Streaming

复杂度分析：从 FLOPS 到 HBM IOs 传统计算的优化，比如 gemm 的 cublas 实现分析都是专注于减少一个计算中 gpu flops 的数量，但是实际上 gpu 的高速内存 SRAM 和相对低速内存 HBM 的读写 IO 率也是影响整体程序性能的重要的因素；因此本文提出做一个 IO-aware attention kernel 来减少 IO 吞吐数量 Tiling 传统 gemm 中的 tiling 是指每次处理一个 block 而不是完整读写一整个行向量，这样可以提升 io 的连续性从而提升性能；但是这里的 attention 计算不能直接用 gemm 的优化，因为存在一个 需要全局 reduce 的 softmax 函数，并且需要用到全局最大值来做 max trick，这个其实是很不利于 tiling 这样分块的实现的 softmax 的背景是 softmax(QKT)softmax(QK^T)softmax(QKT), 对应的矩阵表示的是 query 与 key 之间的相似度/匹配度，然后权重的计算其实是对每一个 query 沿着不同的 key 去做的 red ...

Challenge GPU 并不像 CPU 那么适合做 branching 操作，如果我们在 kernel 中写一个 if-else 指令，那么两个 branch 的函数执行都会被放到解释器这里来，导致执行速度慢于完全并行同等操作; 如果让 batch inference 请求同时通过这个 moe-FFN 层, 这个 batch 内必须统一 expert 而不能分开; 主要框架 标准公式 y=Σi=1nGi(x)Ei(x)y = \Sigma_{i=1}^{n} G_i(x) E_i(x) y=Σi=1n​Gi​(x)Ei​(x) 这里的 GGG 指的是 gating, 用于将输入的 x 向量映射到 n 个 experts 对应的输出用于表示权重; 这里 G 就是一个可训练的神经网络, 可以写作 G=softmax(x⋅Wg)G = softmax(x \cdot W_g) G=softmax(x⋅Wg​) 但是我们用 moe 其实是希望减少计算量，所以需要让很多 expert 模块完全不运行, 这里就需要用类似 top-p 的算法来只要 cumulative 前 90%90\ ...