本文深度解析H100 GPU上高性能矩阵乘法内核的设计原理，涵盖内存层次、张量核心、TMA、数据交织、异步流水线及希尔伯特调度等核心技术。
在GPU的世界里，矩阵乘法就是AI的命脉！你可能不知道，Transformer模型90%以上的计算量，都花在了矩阵乘法上——不管是训练还是推理，它都是那个默默扛起整个大模型运转的“苦力”。
真正决定大模型能不能跑得快、省不省电、能不能落地的核心技术：高性能矩阵乘法内核（matmul kernel）！这篇文章，带你一层一层扒开英伟达H100 GPU的“内脏”，看看顶级工程师是怎么榨干每一块晶体管、每一纳秒、每一度电，把性能推到物理极限的！
首先，作者亚历克萨·戈尔迪奇（Aleksa Gordić）是深度学习系统领域的硬核实战派，长期专注于GPU底层优化、高性能计算和大模型推理加速。他不仅亲手在H100上跑通了接近甚至超越cuBLAS（英伟达官方数学库）性能的自研内核，还把整个过程像解剖青蛙一样，从硬件架构、内存模型、指令集、编译器行为，一直讲到调度策略和空间填充曲线，逻辑严密、细节爆炸。他的博客被无数AI基础设施工程师奉为“圣经”，连PyTorch核心开发者都主动帮他审稿。可以说，他是站在GPU性能金字塔尖上跳舞的人。
那么，为什么矩阵乘法这么重要？因为所有神经网络的线性层、注意力机制里的QKV投影、输出层，本质上都是矩阵乘法。而GPU之所以能成为AI时代的“核武器”，就是因为它的架构天生适合做这种高度并行的计算。
但“适合”不等于“自动高效”——如果你写的内核像教科书那样naive（天真），性能可能只有理论峰值的1%！比如作者举了个例子：只是把一行代码里的“%”和“/”互换了一下，性能就从3171 GFLOP/s暴跌到243 GFLOP/s，整整13倍！为什么？因为破坏了全局内存（GMEM）的访问连续性，触发了DRAM的行激活惩罚。这就是硬件物理规律给你的耳光。
要写出高性能内核，你必须对GPU的内存层次结构了如指掌。H100的内存就像一座金字塔：最底层是容量大但速度慢的显存（VRAM），中间是L2缓存，再往上是每个流式多处理器（SM）私有的L1缓存和共享内存（SMEM），最顶层是寄存器（RMEM）。数据越靠近计算单元，带宽越高、延迟越低。高手的策略就是：把最频繁访问的数据，比如矩阵的小块（tile），提前搬到SMEM甚至寄存器里，尽量减少访问显存的次数。而Hopper架构新增的“张量内存加速器”（TMA），更是让这个过程变得异步、高效、还能自动解决共享内存的“银行冲突”问题。
说到共享内存，这里有个魔鬼细节：它被分成32个“银行”（bank），每个银行一次只能服务一个地址。如果一个线程束（warp，32个线程）里的多个线程同时访问同一个银行的不同地址，就会发生“银行冲突”，请求被串行化，性能直接打骨折。为了解决这个问题，Hopper引入了“数据交织”（swizzling）技术。简单说，就是通过一个精巧的XOR位运算，把原本会冲突的内存地址，重新映射到不同的银行上。作者甚至手把手推导了128B交织模式的掩码生成公式，让你明白为什么无论是按行还是按列访问，都能实现零冲突的单周期加载。
当然，光有内存优化还不够，真正的性能核弹是“张量核心”（Tensor Core）。这是英伟达专门为矩阵运算设计的硬件单元。在Hopper上，它通过一种叫wgmma.mma_async的异步指令，让4个线程束（共128个线程）组成一个“线程束组”（warp group），协同完成一个64x64x16的矩阵乘加运算。作者展示了如何用几行内联PTX汇编，就调用这个强大的硬件单元，把原本需要成百上千行CUDA C++代码才能实现的“线程分块”、“寄存器分片”等复杂逻辑，全部交给硬件自动处理。
但故事还没完。顶级内核的精髓在于“流水线”。作者描述了一种“生产者-消费者”模型：一个线程束组专门负责用TMA从显存预取下一批数据到共享内存的环形缓冲区里，而另一个（或多个）线程束组则专心用张量核心计算已经就绪的数据。两者通过共享内存里的“异步屏障”（barrier）进行同步，实现了计算和数据搬运的完美重叠。这样一来，张量核心和TMA引擎就永远不会闲着，GPU的“光速性能”（Speed of Light）被真正压榨出来。
更绝的是调度策略。传统的做法是给每个输出矩阵块分配一个线程块，顺序执行。但作者引入了“希尔伯特曲线”（Hilbert Curve）这种空间填充曲线来安排计算顺序。为什么？因为它能最大化数据的局部性，让相邻的计算任务尽可能访问相邻的内存区域，从而减少L2缓存和显存的访问次数。在多GPU或多SM协同的场景下，这种调度策略带来的性能提升是惊人的。
最后，作者还提到了“持久化内核”（Persistent Kernel）这种高级技巧。它不是为每个任务都启动一次内核，而是让少数几个内核常驻在SM上，自己从一个任务队列里不断拉取新活儿干。这样不仅省去了内核启动的开销，还能把输出结果的写回操作，和下一批输入数据的加载操作重叠起来，进一步隐藏延迟。
从最朴素的三重循环，到利用TMA、张量核心、异步流水线、希尔伯特调度的SOTA（State-of-the-Art）内核，作者一步步展示了性能是如何从32 TFLOP/s飙升到764 TFLOP/s的。这背后不是魔法，而是对硬件物理极限的深刻理解和极致尊重。正如他所说：“计算机是可以被理解的。”只要你愿意沉下心，一层一层地去解构它。