阿甘的博客 | Arganzheng's Blog

通信与互联（07）：推理侧的通信——custom all-reduce 与 KV 传输

Communication on the Inference Side: Custom All-Reduce and KV Cache Transfer

前六篇建立了一条完整的路径：第一篇的 α-β 模型给出任何一次通信的理论下界，第二、三篇给 α 和 β 填上 NVLink、PCIe、InfiniBand 与 RDMA 的真实数字，第四篇讲 NCCL 如何把这些硬件能力组织成一次 ncclAllReduce，第五篇讲 PyTorch 如何在 stream 上使用它，第六篇把这一切变成 nccl-tests 的曲线和一棵排障决策树。这些内容的默认场景是训练：消息几十 MB 到几 GB、参与者固定、通信可以和反向计算重叠。推理把这套方法推到两个极端。第一个极端是张量并行的 decode：batch 很小、hidden 维度固定，每一层两次 all_reduce，每次只有几十到几百 KB，一个 decode step...

Posted by arganzheng | June 24, 2026

通信与互联（06）：nccl-tests、调优与排障——从带宽曲线到 hang

nccl-tests, Tuning and Debugging Hangs: From Bandwidth Curves to Flight Recorder

前五篇建立了一条完整的因果链：第一篇的 α-β 模型给出一次集合通信的理论时间，第二、三篇给出链路能提供的 β 和 α，第四篇讲 NCCL 如何在探测到的拓扑上选 ring/tree、channel 数、算法与协议去逼近这个上限，第五篇讲 ProcessGroupNCCL 如何把 NCCL kernel 放到自己的 stream 上、watchdog 如何盯着每一个 WorkNCCL。这条链上每一环都可能出错，而错误的表现只有三种：慢、卡、结果不对。本篇的任务是把前五篇变成一套可操作的方法——面对这三种现象，先测什么、先看什么、先改什么。性能这一半的核心工具是 nccl-tests。它做的事情很简单：对一组消息大小依次调用同一个集合通信原语，报告时间、algbw...

Posted by arganzheng | June 20, 2026

通信与互联（05）：PyTorch 的通信栈——ProcessGroupNCCL、stream 语义与计算通信重叠

PyTorch's Communication Stack: ProcessGroupNCCL, Stream Semantics, and Compute-Communication Overlap

上一篇沿着 ncclAllReduce 走完了 NCCL 内部的全部路径：bootstrap、拓扑探测、ring/tree 搜索、transport 建连、调优表、enqueue、一个 kernel 里 nChannels 个 block 各跑一条环，跨机时由 proxy 线程替 GPU 驱动网卡。那一篇的结论可以压缩成一句话：NCCL 是一个把集合通信编译成 CUDA kernel 的库，ncclAllReduce(sendbuf, recvbuf, count, dtype, op, comm, stream) 的最后一个参数决定了这个 kernel 排进哪条队列。这一篇就从这最后一个参数开始。绝大多数人不直接调 NCCL，而是写 dist.all_redu...

Posted by arganzheng | June 17, 2026

通信与互联（04）：NCCL 架构——拓扑探测、channel、算法与协议

NCCL Architecture: Topology Detection, Channels, Algorithms and Protocols

前三篇把”硬件能做到什么”的上限摆出来了：第一篇给出 α-β 模型和 ring all_reduce 的 \(T_{\text{ring}} = 2(n-1)\,\alpha + \frac{2(n-1)}{n}\cdot\frac{S}{\beta}\)；第二篇给 α 和 β 填上数字——NVLink 每 GPU 双向合计 600/900 GB/s（A100/H100），PCIe 4.0 x16 单向约 32 GB/s，IB NDR 单向 50 GB/s——并用 nvidia-smi topo -m 的 NV#/PIX/PXB/PHB/SYS 描述任意两个设备之间的路径；第三篇讲清了跨机时网卡如何绕过 CPU 和主机内存直接读写显存。这些都是”能力”。本篇讲的是 ...

Posted by arganzheng | June 13, 2026

通信与互联（03）：RDMA 与 GPUDirect——绕过 CPU 和主机内存的数据通路

RDMA and GPUDirect: Bypassing the CPU and Host Memory

上一篇给出了节点间链路的物理上限：一张 NDR InfiniBand 网卡 400 Gb/s，约 50 GB/s 单向；一张 HDR 200 Gb/s，约 25 GB/s。也给出了它挂在哪里：每张 GPU 配一张网卡，两者在同一个 PCIe switch 下（nvidia-smi topo -m 里的 PIX 或 PXB），走 PCIe 4.0 x16（单向约 32 GB/s）或 PCIe 5.0 x16（单向约 64 GB/s）。这些数字是链路的能力，不是软件能拿到的带宽。软件拿到多少，取决于数据从显存到网线之间经过了什么。 TCP/IP 的路径上有内核协议栈、至少两次内存拷贝和一个必须参与每个包收发的 CPU。这条路径在 10 Gb/s 时代设计得很合理，到 ...

Posted by arganzheng | June 10, 2026

通信与互联（02）：硬件互联——PCIe、NVLink、NVSwitch 与网络拓扑

Hardware Interconnect: PCIe, NVLink, NVSwitch and Network Topology

上一篇把一次通信的时间写成 \(T = \alpha + S/\beta\)：α 是固定开销，β 是带宽，S 是字节数；ring all_reduce 在 n 个参与者上的时间是 \(T_{\text{ring}} = 2(n-1)\,\alpha + \frac{2(n-1)}{n}\cdot\frac{S}{\beta}\)。这个模型能算出”8 卡 1 GB 的 all_reduce 在 25 GB/s 的链路上大约 70 ms”，但它有两个空位：α 和 β 是多少，取决于数据走的是哪条链路。同一台机器上，两张 GPU 之间的 β 可能是 450 GB/s，也可能是 25 GB/s，也可能只有 10 GB/s，差 40 倍；α 可能是 2 微秒，也可能是 20 ...

Posted by arganzheng | June 6, 2026

通信与互联（01）：集合通信原语与代价模型——α-β 模型与 ring all-reduce

Collective Communication Primitives and the Alpha-Beta Cost Model: Deriving Ring All-Reduce

总纲把这个系列要回答的问题定为一句话：一次 all_reduce 从调用到完成，数据在 PCIe、NVLink、InfiniBand 上是怎么流动的，为什么有时候是带宽的问题、有时候是延迟的问题。要回答它，先得有一把尺子。没有尺子，nccl-tests 打出来的 busbw 23.1 GB/s 只是一个数字，profiler 里 ncclDevKernel_AllReduce 的 145 µs 也只是一个数字，你不知道它们是好是坏、离上限多远、差的那部分该去哪一层找。这一篇造这把尺子。它不碰任何硬件、不读 NCCL 的源码，只做三件事：把集合通信的原语定义清楚，把描述一条链路的 α-β 模型建立起来，再把 ring 和 tree 两个最基本的 all_reduc...

Posted by arganzheng | June 3, 2026

通信与互联：从 NCCL 到 RDMA（总纲）

Communication and Interconnect for AI-Infra, from NCCL to RDMA

内容简介《通信与互联：从 NCCL 到 RDMA》是一组共八篇的系列文章，面向已经跑过多卡训练或多卡推理、准备向下进入通信层的工程师，系统讲解 GPU 之间的数据是怎么流动的：经过哪些硬件链路、由哪些软件层驱动、每一层的代价是什么、出问题时到哪里去看。它回答的问题是：一次 all_reduce 从调用到完成，数据在 PCIe、NVLink、InfiniBand 上是怎么流动的？为什么有时候是带宽的问题，有时候是延迟的问题？1 站在框架的层面看，通信是一行代码：dist.all_reduce(t)。它返回得很快，然后在某个时刻”完成”。Profiler 里它是一段 ncclDevKernel_AllReduce_* 的时间条；nccl-tests...

Posted by arganzheng | June 1, 2026

LoRA 专题（04）：系列总结与通关自测

LoRA Series: Recap and Final Self-Test

三篇正文回答了一个问题：为什么给一个几十亿参数的模型加两个瘦矩阵就能把它微调好，每个旋钮该取多少，训完的东西怎么上线。第一篇讲低秩假设、\(W + \frac{\alpha}{r}BA\) 的梯度与四本账，并直接检验全量微调的 \(\Delta W\) 有多低秩；第二篇把十三种配置在同一份数据上各训 80 步，每个旋钮给一组数字；第三篇从 adapter 文件走到 multi-LoRA 服务。本文把三篇压成一张速查表，拎出贯穿三篇的四条线，列常见误区，再给一套自测。一、总览一张速查表：三篇正文的速查表问题答案在哪篇 LoRA 的公式 \(...

Posted by arganzheng | May 29, 2026

LoRA 专题（03）：工程：adapter 文件、合并、多 LoRA 服务与参考模型

LoRA 03: In Production — Adapter Files, Merging, Multi-LoRA Serving and the Free Reference Model

更新 @2026-09-30：实验用 peft 0.21.1、transformers 5.17.0、bitsandbytes 0.50.2、PyTorch 2.14（CPU），模型 Qwen2.5-0.5B，adapter 是第二篇训出的 r16_all（\(r = 16\) 全部线性层）与 r16_attn（\(r = 16\) 只挂 attention）。配套脚本 ai-learning-labs/lora/03_deploy.py（子实验 files / merge / quant / multi / tokens）。前两篇回答了”为什么”与”怎么配”，训出来的东西是一个目录：adapter_config.json 加一个几十 MB 的 adap...