主要观点总结
DeepSeek开源的DeepGEMM库实现了超越专家级优化的矩阵乘法,利用Hopper GPU达到了惊人的计算速度。DeepGEMM采用简洁的设计,仅300行代码,但性能可与专家调优的内核相媲美。文章介绍了DeepGEMM的性能测试结果,核心优化策略,以及引起的广泛关注和评价。
关键观点总结
关键观点1: DeepGEMM实现超越专家级优化的矩阵乘法
DeepGEMM利用Hopper GPU的惊人速度实现了高效的矩阵乘法,不仅在性能上超越了专家调优的内核,而且采用了简洁的设计。
关键观点2: DeepGEMM的性能表现
DeepGEMM在性能测试中表现出色,特别是在小批量处理的情况下,加速比高达2.7倍。对于混合专家模型和大批量处理,DeepGEMM也提供了稳定的性能提升。
关键观点3: DeepGEMM的核心优化策略
DeepGEMM采用了多项优化技术,包括使用CUTLASS设计中的线程束专用架构、Warp专用内核、张量内存加速器(TMA)等,以提高计算效率和GPU利用率。
关键观点4: DeepGEMM引起的关注和评价
DeepGEMM的发布引起了广泛的关注和积极的评价,人们普遍认为其在矩阵运算中的性能表现超出预期,特别是在高效性和简洁性方面。一些行业观察人士表示,这是通过系统团队串联上下游进行协同优化的典范。
正文
DeepSeek开源周第三弹来了。仅用300行代码就实现超越专家级优化的矩阵乘法?DeepSeek开源的DeepGEMM做到了,不仅在Hopper GPU上飙出1350 TFLOPS的惊人速度,还实现了教科书般简洁。
DeepGEMM是一个一个支持密集和MoE GEMM的FP8GEMM库,为V3/R1训练和推理提供支持。它的核心亮点包括:
通俗来说,DeepGEMM就像是一个超高效的计算工具,专门用于大模型中最常见的数学运算:矩阵乘法。它的特别之处在于使用了FP8(8位浮点数)格式,这种格式可以大大提高计算速度和内存效率,但通常会损失一些精度。DeepGEMM通过精细的缩放技术解决了精度问题,让计算既快又准。
DeepGEMM完全基于NVIDIA的CUDA并行计算平台编写,充分利用了NVIDIA Hopper架构的最新张量核心进行优化。它采用即时编译(JIT)技术,无需预编译,可在运行时动态编译内核,提高了灵活性和适应性。为了解决 FP8 张量核心计算可能存在的精度问题,DeepGEMM 使用 CUDA 核心进行两级累加,确保了计算结果的准确性。尤为值得一提的是,DeepGEMM 的核心计算函数仅约 300 行代码,设计极为简洁,避免了像 CUTLASS 和 CuTe 那样复杂的模板,大大降低了学习和使用的门槛。
虽然DeepGEMM设计简洁轻量,但它的性能表现可以媲美甚至超过那些由专家调优的复杂库,尤其是在处理各种不同形状的矩阵时,这使它成为学习 Hopper FP8 矩阵乘法和优化技术的理想资源。
DeepSeek的研究人员在搭载NVCC 12.8的H800GPU上,测试了DeepSeek-V3/R1推理中可能用到的所有矩阵形状(包括预填充和解码阶段,但不包括张量并行)。所有加速比都是与基于 CUTLASS 3.6内部精心优化的实现相比较得出的。
需要注意的是,DeepGEMM在某些矩阵形状上表现不是特别理想。DeepSeek也表示欢迎有兴趣的开发者提交优化的PR。
性能测试报告展示了DeepGEMM与现有技术相比的性能优势。DeepGEMM 在多数项目中都获得了不错的名次。
在小批量(M=64或128)的情况下,DeepGEMM性能表现尤为出色,加速比高达 2.7 倍。这类似于在短跑比赛中,DeepGEMM展现出了显著的速度优势。这对于AI模型的实时推理(如聊天机器人生成回复)特别有价值。
对于混合专家模型(MoE)的计算,DeepGEMM提供了约1.2倍的稳定性能提升。这种稳定的性能提升对于模型的整体效率非常重要。
在处理大批量数据时,DeepGEMM同样保持了约1.1-1.2 倍的性能优势。这对于批量处理大量文本或图像非常有用,就像工厂的流水线能够更快速地处理大量产品一样。
总体来看,DeepGEMM在小批量处理上的表现特别优异(加速比达到2.7 倍),而在大型矩阵和混合专家模型上也保持了稳定的性能优势(加速比约1.1-1.2倍)。这使得它在各种 AI 模型推理场景中都具有实用价值,尤其是在需要快速响应的应用中。
DeepGEMM 的内核采用了 CUTLASS 设计中的线程束专用架构,这使得数据移动、张量核心 MMA 指令和 CUDA 核心提升可以重叠执行。下图简单展示了这个过程:
为什么这样设计很厉害?传统方式是:先搬运数据,等完全搬完后才开始计算,计算完一批后再搬运下一批。这样 GPU 的大部分部件经常处于等待状态。而 DeepGEMM 的方式是:当搬运工在搬运新数据时,计算工人已经在处理先前搬来的数据了。不同组的计算工人还会交替工作,确保 GPU 的计算单元始终忙碌。
-
Warp 专用内核:基于 CUTLASS 设计,实现数据移动、张量核心 MMA(矩阵乘加)指令及 CUDA 核心提升的并发执行,以提升吞吐量。
-
张量内存加速器(TMA):利用 Hopper 架构中的 TMA 实现异步、高速数据传输,包括 LHS/RHS 矩阵加载、输出存储、LHS 的多播以及描述符预取。
-
专用 PTX 指令:采用 stmatrix 实现高效的线程束级别矩阵存储,并针对线程组进行寄存器数量控制,以优化资源分配。
-
重叠操作:最大化 TMA 存储与非 TMA 右操作数缩放因子加载的重叠,此技术在 CUTLASS 中未见应用
-
统一调度器与光栅化:采用单一调度器处理所有内核类型及线程块光栅化,以提升 L2 缓存复用率。
-
即时(JIT)编译:基于全 JIT 的设计在运行时编译内核,将 GEMM 形状、块大小和流水线阶段视为常量以节省寄存器并进行编译器优化,同时完全展开 MMA 流水线。
-
未对齐的块大小:支持非 2 的幂次方的块大小(如 112),以最大化流式多处理器(SM)利用率,适应不规则形状,提升可扩展性。
-
SASS 级微调:通过翻转编译二进制文件中 FFMA(融合乘加)指令的 yield 和 reuse 位,提升 warp 级并行性和指令重叠,适用于浮点运算