反向 A Kernel 优化全解析（小白版，从零讲起）
日期：2026-07-21
GPU：H20-3e（sm_90a，Hopper）。
本文对象 = 反向传播的 Kernel A（串行 dstate 反向扫描），当前冠军源码
csrc/smxx/bwd_a_dstate_scan_s12.cuh（= 优化日志 docs/20260715-bwd-kernelA-optimization-...
反向 B Kernel 优化全解析（小白版，从零讲起）
日期：2026-07-20
GPU：H20-3e（sm_90a，Hopper）。
本文对象 = 反向传播的 Kernel B（split 成 B1/B2/B3/B0 四个），当前冠军源码
csrc/smxx/bwd_b_split_fuse_b2_b3_sinfix_a1d.cuh（= 文档 l8 + 两项文档未记录的补充
sinfi...
DeepSeek-V4 的 CSA + HCA 详解
说明
本文不基于 DeepSeek-V4 正式论文（作者手上没有该论文）。所有"公式"都是从 sglang 源码的计算语义反推、并逐条标注 文件:行号，可对照代码核验。
公式用 LaTeX 书写（在支持数学渲染的 Markdown 预览里查看）。
记号： 点积， 逐元素乘， 对 求和。
姊妹文档：docs/mla_explaine...
ReplaySSM —— GDN/KDA 线性注意力的 buffered output-only decode
本文是对上游 PR [GDN][KDA] ReplaySSM buffered output-only decode for LinearAttention（sgl-project/sglang #28451，作者 yuan-luo，base main，13 文件+1722/-...
PR #28695 详解：GDN ReplaySSM Ring Spec-Verify
本文件逐字记录了关于 SGLang PR
#28695
「[GDN] Support ReplaySSM Ring Spec-Verify」的全部讲解。
它是 #28511 的
Part B（移植 Dao AI Lab 2026 的 ReplaySSM）。Part A（decode 版）见
姊妹文档 ...
链接：Qwen
tilelang_fused_chunk_gdr_fwd_kernel 流水线机制详解
本文面向第一次接触这个 kernel 的读者，目标不是逐行翻译代码，而是帮助你建立一个清晰的整体心智模型：
这个 kernel 在算什么
为什么它要拆成多组线程
为什么要做双缓冲
为什么会有这么多 barrier
一个 chunk 在 kernel 里到底经历了哪些阶段
文中默认讨论...
幸有金雨一场，替青春加冕
🏆Team Falcons · IEM Cologne 2026 Champions
灭蜂斩龙，问鼎科隆！
十载饮冰，难凉热血！
Nikola “NiKo” Kovač，a living legend.
The Crownless King finally bled his way to the Iron Throne.
Veni！Vidi！Vici！
201...
前置知识：KDA算法公式、triton、cuda、c++
K1_bwd “prepare” kernel 逐行精讲
对应源码：csrc/smxx/bwd_kernel1.cuh 的 _flash_kda_bwd_prepare
目标读者：没接触过自动微分 / CUDA kernel 的人。读完应能自己推出 L2 归一化反向、数值稳定的 、以及门控的逆向 cumsum，并理解代码为何这样写...
前置知识：KDA算法公式、triton、cuda、c++
K2_bwd 反向递推 kernel 逐行精讲
对应源码：csrc/smxx/bwd_kernel2.cuh 的 _flash_kda_bwd_recurrence
目标读者：没接触过自动微分 / CUDA kernel 的人。读完应能自己把每一条梯度公式推一遍，并理解代码为什么这么写。
约定：d<量> 表示 loss...
前置阅读
reed博士：(4 封私信 / 2 条消息) cute 之 Hopper MBarrier - 知乎
竹佬：(4 封私信 / 2 条消息) 写给大家看的 CuTe 教程：async pipeline - 知乎
12 Async Pipeline
对应代码：12_async_pipeline.cu
需要 GPU（SM90 / Hopper）。
核心概念
Pipe...
