阿甘的博客 | Arganzheng's Blog

通信与互联（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 上是怎么流动的？为什么有时候是带宽的问题，有时候是延迟的问题？站在框架的层面看，通信是一行代码：dist.all_reduce(t)。它返回得很快，然后在某个时刻”完成”。Profiler 里它是一段 ncclDevKernel_AllReduce_* 的时间条；nccl-tests ...

Posted by arganzheng | June 1, 2026

GPU Kernel 工程（11）：系列总结与通关自测

GPU Kernel Engineering: Series Recap and Final Self-Test

十篇正文回答了一个问题：一个 kernel 为什么快、为什么慢，以及如何把它写到接近硬件极限。第一篇把 GPU 拆开并建立 Roofline，第二篇写出第一个 kernel 并学会测量，第三、四篇把 memory-bound 的 elementwise 与 reduction 推到带宽墙，第五、六篇把 GEMM 从 naive 推到 Tensor Core，第七篇用 Triton 看编译器接管了哪一层，第八、九篇把这些工具用到 attention、量化与融合 kernel 上组装出一个 decoder layer，第十篇讲怎么剖析、测试、接入框架并合入一个 PR。本文不讲新内容，做三件事：把十篇压成一张表与十段回顾，把贯穿全系列的几条线拎出来，然后给一套三段式的...

Posted by arganzheng | May 20, 2026

GPU Kernel 工程（10）：剖析、测试与贡献——把 kernel 做成产品

Profiling, Testing and Contributing: Turning a Kernel into a Product

前九篇结束时，手上有一个用自己写的 kernel 跑通的 decoder layer 前向：RMSNorm、RoPE、BF16 Tensor Core GEMM、FlashAttention 前向、SiLU-mul、fused residual+RMSNorm、INT4 weight-only GEMM。它们能跑、结果和 PyTorch eager 对得上、每一个都在自己的 benchmark 里比 naive 版本快很多。但”能跑”和”能合入”之间还有一整段工程。一个 kernel 要成为别人敢用的东西，需要回答四个问题：它到底卡在哪里（剖析）；它在所有会遇到的输入上都对（测试）；它确实比原来快、而且以后不会悄悄变慢（benchmark）；它在别人的 GPU ...

Posted by arganzheng | May 20, 2026

GPU Kernel 工程（09）：量化与融合 kernel——推理系统的其余部分

Quantized and Fused Kernels: The Rest of the Inference Stack

上一篇把 attention 讨论完了。一个 decoder layer 里除了 attention 和标准 GEMM，剩下的是一堆”小 kernel”：RMSNorm、RoPE、SiLU-mul、把 KV 写进分页 cache、把权重从 INT4 解开、把激活压成 FP8、MoE 的 token 重排、采样。它们单个都不复杂，但数量多、变化快，加起来占掉推理时间的一个可观比例——而且是 vLLM、SGLang 这些项目里 PR 最活跃的区域。这一篇把它们放在同一个方法论下过一遍：先算这个 kernel 理论上要搬多少字节、做多少 FLOPs，再看实现，再解释差距。核心问题是总纲提出的那个：一个 INT4 weight-only GEMM，decode...

Posted by arganzheng | May 19, 2026

GPU Kernel 工程（08）：Attention Kernel——FlashAttention 与 PagedAttention

Attention Kernels: FlashAttention and PagedAttention from Derivation to Code

前七篇分别处理了 GPU 的硬件结构与 Roofline、CUDA 执行模型、访存合并、shared memory 与 reduction（其中包括 online softmax）、GEMM 的分块、Tensor Core 与 CUTLASS、Triton。这一篇是它们的汇合点：attention 同时包含两个 GEMM（\(QK^T\) 与 \(PV\)）、一个逐行的 reduction（softmax），以及推理时特有的内存访问模式（分页的 KV cache）。它是 Transformer 推理里最重要、也最难写好的 kernel。总纲给这一篇提的核心问题是：一个序列长度 4k、head dim 128 的 attention，标准实现和 Flas...

Posted by arganzheng | May 18, 2026

GPU Kernel 工程（07）：Triton——块级编程与编译器的边界

Triton: Block-Level Programming and Where the Compiler Stops

前六篇一直在 CUDA 的世界里：每个线程算什么、warp 怎么合并访存、shared memory 怎么分块、mma.sync 怎么喂 fragment。到了第六篇，一个能跑到 cuBLAS 七八成性能的 BF16 GEMM 已经是两三百行代码，而且每一行都有”为什么这样写”的理由——tile 尺寸、bank conflict 的 padding、cp.async 的 stage 数、寄存器分块的形状。这一篇换一种写法。Triton 把”线程”从编程模型中拿掉：程序员以 block 为单位思考，写的是”这个 block 加载哪一块数据、做什么张量运算、存到哪里”，而线程到元素的映射、shared memory 的分配、向量化、软件流水、mma 指令的选择，全部...

Posted by arganzheng | May 17, 2026

GPU Kernel 工程（06）：Tensor Core、CUTLASS 与 CuTe

Tensor Cores, CUTLASS and CuTe: Programming the Matrix Units

上一篇用 CUDA Core 把 GEMM 的分块结构讲透了。回顾一下那个结构，因为本篇要做的事情就是把它”接”到另一种计算单元上：一个 thread block 负责输出矩阵 \(C\) 的一个 \(BM \times BN\) 的 tile，沿 \(K\) 维以 \(BK\) 为步长循环；每一步把 \(A\) 的 \(BM \times BK\) 子块和 \(B\) 的 \(BK \times BN\) 子块搬进 shared memory，全局读取量从 naive 的 \(2MNK\) 个元素降为 \(MNK \cdot (1/BM + 1/BN)\)，\(BM = BN = 128\) 时减少 64 倍；每个线程再从 shared me...

Posted by arganzheng | May 16, 2026

Arganzheng's Blog