开源第三弹DeepSeek,完美榨干GPU,FP8训练秘籍公开。

智慧东西2月26日报道,刚刚,DeepSeek开源周第三弹发布。——DeepGEMM,一个支持密集和MoE FP8GEMM 为V3/R1练习和推理增加GEMM库的动力。
- ⚡ Hopper GPU性能高达1350 FP8 TFLOPS
- ✅ 依赖性不大,像教程一样干净。
- ✅ 完全即时编译JIT(无需预编译安装)
- ✅ 简化设计:核心逻辑约为300行- Kernels在大多数矩阵大小上都比专家调整的要好。
- ✅ 支持密集(Dense)布局和MoE的两种布局

GitHub:
眼尖的网友已经在工程推广人名单上捕捉到了一个“Liang”,并在DeepSeek文章评论区提问:“是梁文锋(DeepSeek创始人)吗?

DeepGEMM是一个FP8通用矩阵乘法矩阵乘法矩阵,专门用于清洁和高效。(GEMM)而且设计的库具有DeepSeek-V3等粗粒度扩展功能。它支持一般和混合专家。(MoE)分组GEMM。该库由CUDA编写,安装过程中不需要编译,而是使用轻量级即时(JIT)所有的kernel都是在模块运行时编译的。
根据DeepSeek曝光的数据,一般GEMM(密集模型)中矩阵运算可以加速高达2.7倍,在GEMM(MoE模型)中分组连续布局,在掩码布局下可以加速高达1.2倍。
目前,DeepGEMM只支持英伟达Hopper Tensor Core。FP8解决不准确的问题 Tensor Core累计问题,采用CUDA核心二级累计(改进)。
虽然它使用了CUTLASS和CuTe的一些概念,但是它避免了对其模板或代数的过度依赖。取而代之的是,该库的设计非常简单,只有一个核心kernel函数,包括大约300行代码。它使它成为学习Hopper 干净易访问的FP8矩阵乘法和优化技术资源。
尽管DeepGEMM设计轻巧,但其性能与各种矩阵形状的专家调节库相当或超过后者。
配备NVCCCCC的DeepSeek DeepSeek-V3/R1推理中使用的所有形状(包括预填充和解码,但没有张量并行性)都在12.8H800上进行了测试。所有加速指标都是基于CUTLASS。 对比计算3.6内部精心优化的实现。
DeepGEMM在某些形状上的表现并不好,所以DeepSeek欢迎开发者对PR进行优化。矩阵运算的最高加速度是常规GEMM(密集模型)的2.7倍。

连续布局、掩码布局下的GEMM(MoE模型)速度可提高1.1倍~1.2倍。

DeepGEMM一发布,DeepSeek的推文评论区就广受好评。有些人担心英伟达股票:


有些人热情地称赞新代码库和DeepSeek工程师:






DeepSeek分享了一个清晰的上手指南,需要HopperGPU架构。、必须支持sm_90a,要求是Python 3.8、CUDA 12.3、PyTorch 2.1、CUTLASS 或者更新版本3.6。DeepSeek强烈推荐CUDA 为了获得最佳性能,12.8或更高版本。

开发:

安装:

在Python项目中导入deep_gemm,然后就可以开始享受了。
该代码库只包含GEMMMM。 kernel。TMA对齐和转置需要LHS扩展因素,仅支持NT格式(非转置LHS和转置RHS)。对转换或其它FP8转换操作,需要独立实现或将其融入到以前的kernel中。尽管这个库提供了一些简单的PyTorch实用函数,但是这些函数可能会导致性能下降。提高GEMMMMMM的关键在于DeepSeek。 自己的kernels。
该代码库除kernel外,还提供了一些实用的函数和环境变量。
使用DeepSeek?表示排除在CUTLASS之外的技术。根据CUTLASS的设计,DeepGEMM中的核心经历了warp的特殊化,可以实现重叠数据移动、张量核心MMA指令和CUDA核心提升。下面的图表显示了这个过程的简化图表:

1、Hopper TMA功能
张量内存加速器(TMA)它是引入Hopper架构的一种新型硬件功能,旨在实现数据移动速度更快、异步更快。具体而言,DeepSeek使用TMA来实现以下目标:
- LHS、TMA负荷的LHS扩展因子和RHS矩阵
- 导出矩阵的TMA存储
- TMA multicast组播(LHS矩阵特有)
- 预取TMA描述符
2、常见的细节提升
- 使用stmatrix PTX指令
- 对于不同的warpgroups定制的存储器记数控制
- 尽量重叠,例如重叠TMA存储和非TMA存储。 载入RHS扩展因素?
3、块调度器的统一优化
- 适合所有非分组和分组核心的调度程序
- 加强L2缓存器的光栅化
4、全JIT设计?
选择完全即时编译的DeepGEMM。(JIT)在安装过程中,设计不需要编译。所有的核心都是在运行过程中使用轻量级JIT进行编译。该方法具有以下优点:
- 在编译过程中,GEMM的形状、块大小和管道阶段被视为常量。
- 储存寄存器
- 编译器可能会做更多的改进。
- 自动化选择块大小、warpgroups数量、最佳流程阶段和TMA集群大小。
- 但是,如果没有自动调整,最佳方案将被确定地选择。
- 全面开展MMA流程,为编译器提供更多提升机会
- 对小形状至关重要
- 详情请参阅launch__k_iterations kernel文档
总的来说,JIT显著提高了小外观的性能,类似于Triton编译器的方法。
5、块大小不对齐
对某些形状来说,与2的幂对齐的块大小可能会导致SM没有得到充分利用。例如,对M=256, N=BLOCK_典型的块大小分配会7168M=128, BLOCK_N=128造成只有(256 / 128) * (7168 / 128) = 使用112132个SM。
对此,DeepSeek支持未对齐的块大小(例如 112),使(256 / 128) * (7168 / 112) = 在这一场景中,128SM可以工作。当粗粒度扩大时,这项技术的实施需要仔细改进,但最终可以提高性能。
本文来自微信微信官方账号的“智东西”(ID:zhidxcom),作者:ZeR0,编辑:漠影,36氪经授权发布。
本文仅代表作者观点,版权归原创者所有,如需转载请在文中注明来源及作者名字。
免责声明:本文系转载编辑文章,仅作分享之用。如分享内容、图片侵犯到您的版权或非授权发布,请及时与我们联系进行审核处理或删除,您可以发送材料至邮箱:service@tojoy.com




