KKKZOZ’s Blog

本文以 transformers==5.14.1 中的 Qwen3_5TextModel 和 Qwen3_5ForCausalLM 实现为依据，讨论文本主干，不展开视觉编码器和多模态输入。源码参考：modular_qwen3_5.py、cache_utils.py。 总览 Qwen3.5 是由线性注意力层和标准全注意力层交替组成的 hybrid decoder-only Transformer。 文本主干的计算顺序是：Embedding → 多个 Qwen3_5DecoderLayer → 最终 Qwen3_5RMSNorm → lm_head。每个 decoder layer 都包含一个 token mixer 和一个 MLP；token mixer 的具体类型由 config.layer_types[layer_idx] 决定。 在默认配置中，full_attention_interval=4，因此每第 4 个 layer 使用 full_attention，其余 layer 使用 linear_attention。 配置定义 Qwen3_5TextConfig 同时定义模型宽度、层数、两类 attention 的维度和层类型。 参数 含义 影响的模块 hidden_size token 表示维度 H embedding、attention 输出、MLP 输入输出 intermediate_size MLP 中间维度 I gate_proj、up_proj、down_proj num_hidden_layers decoder layer 数量 layers num_attention_heads 全注意力的 query head 数 Nq Qwen3_5Attention num_key_value_heads 全注意力的 key/value head 数 Nkv GQA 和 KV cache head_dim 全注意力每个 head 的维度 Q/K/V 和 RoPE linear_key_head_dim 线性 attention 的 key/query head 维度 Dk Gated DeltaNet linear_value_head_dim 线性 attention 的 value head 维度 Dv Gated DeltaNet linear_num_key_heads 线性 attention 的初始 key/query head 数 Gated DeltaNet linear_num_value_heads 线性 attention 的 value head 数 recurrent state linear_conv_kernel_dim 线性 attention 的 causal convolution kernel 大小 conv1d 和 convolution cache layer_types 每层的 linear_attention 或 full_attention decoder layer 分支 layer_types 未提供时，配置代码按 full_attention_interval 生成： ...

Overview 首先介绍一下使用 tilelang 写 kernel 的两个视角： Block 视角 从一个 BLock 的视角写 Kernel, 操作的对象是整个 tile（比如一个 (BLOCK_M, BLOCK_N) 的矩阵片段），并行的粒度和分工方式通过 T.Parallel/T.copy 等声明式原语来表达，不关心具体哪个 threadIdx 做哪个元素。 换句话说，我们只需要声明，具体的线程分配由 tilelang 框架来完成 Thread 视角 从单个 Thread 的视角写 Kernel，block 内的并发是隐含的——因为 BLOCK_N 个 thread 各自都在执行同一份代码，天然并行，但代码里不体现这个并行结构。 Note Thread 视角可以让我们更贴近 CUDA 代码，但好像 tilelang example 中的用法基本都是 Block 视角，所以我们的重心应该放在 Block 视角上 ...

在生产环境下使用 TileLang，需要考虑几个问题： 算子缓存 与 torch.compile 兼容 算子缓存 算子定义 在了解算子缓存之前，我们需要了解一个 GPU 算子应该如何定义，通常由三部分组成： 参数类型 RMSNorm 中的实例 作用 问题规格（Workload / specialization parameters） M, N, eps, dtype 确定一类具体计算问题 调度参数（Schedule / config / meta-parameters） block_m, threads 确定该 workload 的一种实现方案 运行时参数（Runtime arguments / operands） x, weight 每次调用时传入的实际数据 在代码里对应得非常直接： _rms_norm_kernel(M, N, eps, dtype) # workload parameters def _func(block_m, threads): # schedule/meta-parameters def main(x, weight, y): # runtime arguments 这三类参数不是一次性传给同一个函数，而是在 kernel 从定义到执行的过程中逐层绑定： jit_impl = _rms_norm_kernel(M, N, eps, dtype) jit_kernel = jit_impl(block_m=4, threads=128) y = jit_kernel(x, weight) 第一步绑定 workload 参数，得到 TileLang JITImpl 描述一个确定的计算问题，例如处理形状为 (M, N)、数据类型为 dtype 的 RMSNorm，但尚未确定具体的调度方案。 第二步绑定 schedule 参数，得到可以执行的 JITKernel block_m/threads 可以来自默认配置、手动配置或 autotune 不同 schedule 对应同一个 workload 的不同实现版本 最后传入 x/weight 等 runtime arguments，启动已经选定的 JITKernel 运行数据只参与本次调用，不改变前面已经确定的 workload 和 schedule 算子调用流程 一次真实的算子调用从业务 Op 开始，依次确定 workload、选择 schedule、取得编译结果，最后才把 Tensor 作为 runtime arguments 启动 kernel ...

ML Compiler Overview ML Compiler 位于深度学习框架和底层硬件之间，作用是把 TensorFlow、PyTorch、JAX 等框架描述的高层张量计算转换成能在 CPU/GPU/TPU 等硬件上高效执行的程序。 它的典型架构是 Frontend → IR → Optimization → Backend: Frontend 接收框架已经构造或捕获出来的计算图/算子程序（例如 tf.function 得到的 TF Graph、PyTorch Dynamo 得到的 FX Graph），转换成统一的中间表示 IR 中间层进行算子融合、常量折叠、布局变换、内存规划、并行化等与机器学习计算相关的优化 Backend 再根据目标硬件进行 lowering、调度和代码生成 因而 ML compiler 输入通常是带有 tensor shape、dtype、算子及数据依赖关系的计算图或 IR，输出则是面向特定设备的低层 IR、kernel 或可执行程序。 需要特别区分的是：graph capture 不一定属于 compiler 本身——例如 TensorFlow 由 tf.function tracing 捕获图，PyTorch 由 TorchDynamo 捕获图，然后 ML compiler 才接手，对这个图做优化并最终生成高效机器码。 IR Pipeline 下面这套 IR 划分描述的是一个典型 ML Compiler 从高层模型计算逐渐 lowering 到硬件可执行 kernel 的过程。不同编译器实际使用的 IR 名称和数据结构可能不同，但抽象层次通常可以归纳为： ...

vLLM 中同时存在多条与 FX 和 torch.compile 相关的执行路径，它们虽然共享 PyTorch 的 graph infrastructure，但解决的问题并不相同。理解这些路径的边界，是区分“模型结构改写”“Inductor 编译优化”和“CUDA Graph replay”的前提。 进入正文前，可以先区分三类用法： HF Transformers modeling backend：直接使用 fx.Tracer 分析 Hugging Face forward，识别可融合结构并通过 AST 改写模型源码 vLLM 自定义 torch.compile 主路径：由 Dynamo 产生完整 FX GraphModule，再由 vLLM 控制 graph partition、shape specialization、custom passes、cache 和 CUDA Graph integration 普通 @torch.compile 小函数：FX 主要作为 Dynamo、AOTAutograd 和 Inductor 的内部 IR，vLLM 通常不直接处理这些 graph Piecewise CUDA Graph 不是第四条独立的 tracing 路径，而是建立在 vLLM Piecewise Compilation 之上的运行时优化：编译阶段先拆分并编译 graph regions，随后只对兼容区域执行 CUDA Graph capture/replay。 这篇文章的核心问题是：vLLM 如何围绕标准 torch.compile 增加一层 LLM-serving-aware 的 graph processing、compilation policy 和 runtime integration。 ...

核心结论 Piecewise CUDA Graph 的核心思想是：不要求整个 Transformer forward 都满足 CUDA Graph capture 条件，而是将 CUDA Graph 不兼容的算子作为边界，只 capture 其余 CUDA-Graph-safe 区域。 在 vLLM 中，最典型的边界是 Attention。因此，PIECEWISE 模式并不是“把一个 CUDA Graph 切开执行”，而是将模型划分为多个独立的 CUDA Graph，并在它们之间以 eager 模式执行 CUDA-Graph-unsafe 操作： CUDA Graph -> Attention (eager) -> CUDA Graph -> Attention (eager) -> CUDA Graph CUDA Graph 会提前 capture 一组 GPU operations，后续通过 graph replay 一次性提交，从而降低 CPU 逐个启动 kernel 的开销。Transformer 中的 RMSNorm、Linear/GEMM、activation、FFN 和 residual 等计算，在 capture size 确定时通常具有稳定的执行结构；Attention 则需要处理 KV Cache、动态 shape 和运行时 metadata，对 CUDA Graph 的兼容要求更高。 ...

在和 GGML 打交道时不知道已经写了几篇文档了，这篇一定是最后一个 这篇文档主要记录我用 GGML 实现一个简单的 LLM Inference Engine 时遇到的问题 C++ Initialization 在 C++ 中， class 中的成员在构造函数体开始之前就被默认构造了，要控制其行为，需要使用初始化列表 class Member { public: Member() { std::cout << "Member 默认构造\n"; } Member(int x) { std::cout << "Member 带参构造: " << x << "\n"; } }; class MyClass { Member m1; // 成员对象 Member m2; int value; public: // 情况1: 不使用初始化列表 MyClass() { std::cout << "构造函数体开始\n"; value = 10; // 这是赋值，不是初始化！ } // 情况2: 使用初始化列表 MyClass(int v) : m1(1), m2(2), value(v) { std::cout << "构造函数体开始\n"; } }; 情况 1 输出 Member 默认构造 // m1 在构造函数体前被默认构造 Member 默认构造 // m2 在构造函数体前被默认构造 构造函数体开始 情况 2 输出 Member 带参构造: 1 // m1 在构造函数体前初始化 Member 带参构造: 2 // m2 在构造函数体前初始化 构造函数体开始 对于指针来说，初始化就是将其设置为 nullptr, 对于 STL 容器，就是初始化空容器 ...

向量内积 物理意义 内积本质上是一个向量在另一个向量方向上的投影长度，与基准向量长度的乘积。 如果提取特征或求取分量，令 $\mathbf{u}$ 为单位向量（$\|\mathbf{u}\| = 1$），则 $\mathbf{v} \cdot \mathbf{u}$ 直接输出 $\mathbf{v}$ 在 $\mathbf{u}$ 方向上的标量投影。 工程应用： 在信号处理（如傅里叶变换）中，信号与正交基函数的内积，就是在提取该信号在特定频率上的能量分量。在经典力学中，功的计算 $W = \mathbf{F} \cdot \mathbf{d}$ 就是提取力在位移方向上的有效分量并相乘。 由于内积公式中包含 $\cos(\theta)$，它是衡量高维空间中两个向量方向“一致性”或“对齐程度”的线性算子。 当 $\mathbf{a} \cdot \mathbf{b} > 0$ 时，夹角为锐角，两者存在正相关性。 当 $\mathbf{a} \cdot \mathbf{b} = 0$ 时，$\cos(\theta) = 0$，两向量正交（垂直）。在工程上，这意味着两个系统、信号或特征完全独立，互不干涉（即协方差为零）。 当 $\mathbf{a} \cdot \mathbf{b} < 0$ 时，夹角为钝角，存在负相关性。 工程应用： 在机器学习和数据挖掘中，将向量归一化后求内积，即为余弦相似度（Cosine Similarity），常用于衡量文本词向量的语义相似性或推荐系统中用户偏好的匹配度。 Note 向量内积天然满足交换律，即 $\mathbf{a} \cdot \mathbf{b} = \mathbf{b} \cdot \mathbf{a}$，即 $\mathbf{a}$ 投影在 $\mathbf{b}$ 上和 $\mathbf{b}$ 投影在 $\mathbf{a}$ 上的数值是相等的 ...

Notes on L1 and L2 norms, KL divergence, and their applications in machine learning.

记录一些我在使用 cli 时简单整理的一些东西 Setup Docker curl -fsSL https://get.docker.com -o get-docker.sh sudo DOWNLOAD_URL=https://mirrors.ustc.edu.cn/docker-ce sh get-docker.sh sudo groupadd docker # optional sudo usermod -aG docker $USER newgrp docker sudo systemctl start docker apt 替换镜像源： For 24.04: sed -i "s@http://.*archive.ubuntu.com@https://mirrors.aliyun.com/@g" /etc/apt/sources.list.d/ubuntu.sources sed -i "s@http://.*security.ubuntu.com@https://mirrors.aliyun.com/@g" /etc/apt/sources.list.d/ubuntu.sources sed -i "s@http://ports.ubuntu.com@https://mirrors.aliyun.com@g" /etc/apt/sources.list.d/ubuntu.sources For 22.04: sed -i "s@http://.*archive.ubuntu.com@https://mirrors.aliyun.com/@g" /etc/apt/sources.list sed -i "s@http://.*security.ubuntu.com@https://mirrors.aliyun.com/@g" /etc/apt/sources.list Network # 查看网关 ip route # 查看 ip 地址 ip addr # 查看每个进程的网速 sudo apt install nethogs nethogs apt sudo 会为了安全重置环境变量, 加上 -E 参数来保留当前的环境变量 ...

Transformer 是基于全局视角处理序列的，它通过 Self-Attention 一次性让所有 token 互相交互。而 RNN（Recurrent Neural Network，循环神经网络）的本质是基于时间步的顺序迭代。 如果把 Transformer 比作同时看到整段句子的“上帝视角”，RNN 则像是按照从左到右的顺序逐字阅读，并在脑海中维护一个不断更新的“记忆向量”。 RNN 的核心结构是一个循环单元。它在处理序列时，并不是一次性输入整个序列（如 [batch, seq_len, dim]），而是将序列沿着 seq_len 维度切开，在每一个时间步 $t$ 输入一个 token。 为了保留历史上下文，RNN 引入了隐藏状态 $h_t$。 $h_t$ 是一个固定维度的向量，它包含了从时间步 $0$ 到 $t$ 的所有历史信息压缩。 在每一个时间步，RNN 接收两个输入：当前时刻的输入 $x_t$ 和 上一时刻的隐藏状态 $h_{t-1}$。 RNN 会使用同一套权重矩阵（参数共享）来融合这两个输入，生成新的 $h_t$。 将循环单元沿时间维度展开后，可以看到隐藏状态是连接相邻时间步的唯一通道；每个时间步的输入只进入当前单元，而单元参数在整个序列中共享： --- config: theme: neutral --- flowchart LR hp["h_{t-1}"] --> rt["RNN 单元<br/>共享参数"] xt["x_t"] --> rt rt --> ht["h_t"] rt --> yt["y_t（可选）"] ht --> rn["RNN 单元<br/>共享参数"] xn["x_{t+1}"] --> rn rn --> hn["h_{t+1}"] rn --> yn["y_{t+1}（可选）"] 这里的“两个 RNN 单元”并不是两个独立的 module，而是同一个 RNN module 在相邻时间步的两次调用。推理时，for 循环每次传入新的 x_t 和当前的隐藏状态，并计算出下一时刻的 h_t；正是这种从 $h_{t-1}$ 到 $h_t$ 的依赖，使 RNN 必须按顺序计算。 ...

Vanilla RNN 的缺陷：为什么我们需要 LSTM？ 在上一篇中我们提到，RNN 的核心公式是 $h_t = \tanh(W_{xh} x_t + W_{hh} h_{t-1} + b_h)$。这种设计在处理长序列时会遇到两个工程上的致命问题： 梯度消失与梯度爆炸（Vanishing/Exploding Gradients） 在反向传播计算梯度时，误差需要沿着时间步反向传递。这会导致权重矩阵 $W_{hh}$ 被连乘 $t$ 次。根据线性代数原理，如果 $W_{hh}$ 的最大特征值小于 1，连乘后梯度会呈指数级衰减趋近于 0（梯度消失）；如果大于 1，则会指数级放大（梯度爆炸）。梯度消失意味着网络根本无法学习到长距离的依赖关系 信息覆盖（Information Overwrite） RNN 只有一个隐藏状态 $h_t$。在每一个时间步，新的输入 $x_t$ 都会强制与历史信息 $h_{t-1}$ 混合。没有任何机制能够保护早期非常重要但最近没有出现的信息。这就好比一个容量有限的栈，新数据不断涌入，旧数据很快就被冲刷掉了 LSTM 的核心思想：分离状态与引入门控机制 为了解决上述问题，LSTM 将内部状态拆分为两条相互配合的通路，并通过门控信号控制信息如何在通路之间流动。 细胞状态 $c_t$ (Cell State)：位于上方的主干通路，负责携带长期记忆。从 $c_{t-1}$ 到 $c_t$ 的更新主要由逐元素乘法和加法构成，因此为梯度跨时间步传播提供了更顺畅的路径。 隐藏状态 $h_t$ (Hidden State)：位于下方的输出通路，既参与当前时间步的门控计算，也作为当前时间步的输出，并传递给下一时间步。 门控机制 (Gating Mechanism)：将上一时刻的隐藏状态 $h_{t-1}$ 与当前输入 $x_t$ 拼接后，分别计算遗忘门 $f_t$、输入门 $i_t$、输出门 $o_t$ 和候选记忆 $\tilde{c}_t$。其中三个门使用 Sigmoid，将每个维度的控制值限制在 $[0, 1]$；候选记忆使用 $\tanh$ 生成待写入的内容。 一个时间步内部的数据流可以按图中的顺序理解：遗忘门决定保留多少 $c_{t-1}$，输入门决定写入多少候选记忆，二者相加得到新的细胞状态 $c_t$；随后，$c_t$ 经过 $\tanh$ 压缩，并由输出门决定其中多少信息暴露为隐藏状态 $h_t$。 ...