QuACK GEMM内核深度剖析Hopper与Blackwell架构性能对比【免费下载链接】quackA Quirky Assortment of CuTe Kernels项目地址: https://gitcode.com/gh_mirrors/quack15/quackQuACKQuirky Assortment of CuTe Kernels是一个基于CuTe DSL的高性能CUDA内核库专门针对现代GPU架构进行优化。本文将深入剖析QuACK中的GEMM通用矩阵乘法内核对比NVIDIA HopperSM90与BlackwellSM100架构的性能差异帮助开发者理解在不同GPU架构上的优化策略。架构演进从Hopper到BlackwellQuACK的GEMM实现针对不同GPU架构采用了专门的优化策略。Hopper架构SM90和Blackwell架构SM100在硬件特性上有着显著差异QuACK通过不同的内核实现充分利用了各自的优势。Hopper架构SM90的核心特性支持WGMMAWarp Group MMA指令可直接从共享内存读取数据支持TMA多播实现跨CTA的数据广播最大集群大小为4共享内存为228KB使用WGMMA进行矩阵乘加运算Blackwell架构SM100的改进引入TMEMTensor Memory和UMMAUnified MMA指令支持更大的集群大小最大16支持硬件随机舍入引入CLCCluster Launch Control进行持久化内核调度性能对比分析通过QuACK的基准测试我们可以看到两个架构在不同工作负载下的性能表现GEMM内核的关键差异1. 内存访问模式Hopper架构使用TMATensor Memory Accelerator多播技术允许单个TMA加载操作将数据广播到集群中的多个CTA。而SM120RTX 5090虽然支持TMA加载和存储但不支持多播功能这影响了多CTA GEMM的数据共享效率。2. 计算核心差异Hopper使用WGMMA指令直接从SMEM读取A/B矩阵Blackwell使用tcgen05.mma指令将累加器写入TMEMSM120使用warp级MMA指令需要通过ldmatrix从SMEM复制到RMEM3. 集群支持能力Hopper的最大集群大小为4而Blackwell支持最大16的集群大小。在QuACK的实现中这直接影响了几何内核的并行效率。实际性能数据根据QuACK的基准测试结果我们可以看到在大型矩阵运算中MNK4096Hopper架构在fp16精度下达到约1500 TFLOP/sBlackwell架构在相同精度下性能提升15-20%SM120由于缺少TMA多播在多CTA场景下性能受限内存带宽利用率Hopper充分利用228KB共享内存Blackwell优化TMEM使用减少寄存器压力SM120受限于99KB共享内存需要更紧凑的内存布局优化策略对比Hopper架构优化在quack/gemm_sm90.py中Hopper GEMM实现采用了以下策略TMA多播优化减少L2内存流量WGMMA直接访问直接从SMEM读取数据多级流水线重叠计算和内存访问Blackwell架构优化在quack/gemm_sm100.py中Blackwell GEMM实现的关键特性持久化瓦片调度更好地重叠内存加载/存储与MMA操作Warp专业化避免主循环加载和MMA之间的显式流水线TMEM利用专用张量内存区域SM120架构适配在quack/gemm_sm120.py中针对SM120的适配显式SMEM→RMEM复制通过ldmatrix指令集群限制cluster_shape(1,1,1)避免多播依赖共享内存优化适应99KB限制配置参数对比QuACK的GEMM配置在quack/gemm_config.py中提供了针对不同架构的优化参数Hopper典型配置tile_shape_mnk (128, 256, 64) cluster_shape_mnk (2, 1, 1)Blackwell典型配置mma_tiler_mnk (128, 256, 64) cluster_shape_mnk (4, 2, 1) use_2cta_instrs True性能调优建议1. 数据类型选择fp16/bf16两个架构都提供良好支持fp8Blackwell提供更好的fp8支持int8/uint8仅Blackwell原生支持2. 瓦片大小优化HopperM必须是64/128N必须是64/128/256BlackwellM可以是64/128use_2cta_instrsFalse或128/256use_2cta_instrsTrueK维度通常设置为4个MMA指令的倍数3. 集群配置根据问题规模和GPU规格选择集群大小考虑内存对齐要求至少16字节对齐平衡计算密度和内存带宽实际应用场景大语言模型推理在LLM推理中QuACK的GEMM内核表现出色注意力机制softmax和线性层的高效实现前馈网络矩阵乘法的优化执行层归一化rmsnorm的高性能计算科学计算大型稠密矩阵运算批量矩阵乘法高精度数值计算未来展望随着GPU架构的不断演进QuACK团队持续优化内核实现支持新数据类型如MXFP4、NVFP4等低精度格式自适应调度根据硬件特性动态选择最优内核跨架构兼容确保代码在Hopper、Blackwell和未来架构上的可移植性总结QuACK的GEMM内核展示了在不同GPU架构上的精细优化策略。Hopper架构通过WGMMA和TMA多播实现了高效的计算而Blackwell架构则通过TMEM和UMMA指令进一步提升了性能。对于开发者而言理解这些架构差异有助于选择最佳的配置参数从而在实际应用中获得最优性能。无论是进行大规模AI模型训练还是高性能科学计算QuACK都提供了经过充分优化的GEMM实现让开发者能够充分利用现代GPU的硬件特性实现极致的计算性能。【免费下载链接】quackA Quirky Assortment of CuTe Kernels项目地址: https://gitcode.com/gh_mirrors/quack15/quack创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考