大模型训练推理(li)神作(zuo)，又(you)更新了！
主(zhu)流大模型都在用的FlashAttention，刚刚升级(ji)第三代。
时隔一年，FlashAttention-3已经全方位升级(ji)。
训练速度提升1.5-2倍，FP16下(xia)计算吞吐量高达740TFLOPs/s，达理(li)论最大吞吐量75%，更充分(fen)利用计算资源，此前只能做到35%。
FP8下(xia)速度接近1.2PFLOPs/s！
同时误差也进一步减小(xiao)，FP8下(xia)的误差比标准Attention减少2.6倍。
而且这一次，不再(zai)是一作(zuo)Tri Dao单(dan)打独斗，FlashAttention-3直(zhi)接和英伟达、Meta、谷歌等合作(zuo)，针对最强芯片H100专门做优(you)化。
英伟达CUTLASS团队(dui)和cuDNN团队(dui)，都直(zhi)接为该研(yan)究提供支持。
同时和前作(zuo)一样，FlashAttention-3也将开源，PyTorch和Hugging Face中(zhong)都集(ji)成(cheng)。
作(zuo)者之一Vijay Thakkar激动表(biao)示：
曾经在FA2发布时，我就说过这句话。今天，我想再(zai)说一次：
看到CUTLASS和CuTe被用来开让Tensor Core大显身手的新算法，真的泰裤辣。
曾经在FA2发布时，我就说过这句话。今天，我想再(zai)说一次：
看到CUTLASS和CuTe被用来开让Tensor Core大显身手的新算法，真的泰裤辣。
前Stable Diffusion老板(ban)Emad也非常关注这一进展(zhan)，他推测使(shi)用FlashAttention-3，能将4090的FP8计算吞吐量推升到700+TFLOPs。
充分(fen)利用Hopper架构特点
自初代发布以来，FlashAttention已经使(shi)大模型速度提高了4-8倍，但还(hai)有一个遗憾：尚未充分(fen)利用现代 GPU。
针对英伟达H100倍后的Hopper架构新特性，三代进行(xing)了专门优(you)化。
整个系列的核心思路，是IO感知优(you)化和分(fen)块处理(li)。
作(zuo)者认(ren)为，传统的注意力机制效(xiao)率低的原因，在处理(li)长序列时，会出现内存访问(wen)操作(zuo)频繁，以及算法复杂度指数级(ji)暴(bao)增(zeng)这两大问(wen)题(ti)。
FlashAttention通(tong)过IO感知优(you)化将数据从(cong)较大但缓慢的高带宽内存（HBM）加载到较小(xiao)但更快的片上内存（SRAM），在SRAM中(zhong)执行(xing)计算，减少了内存读写(xie)操作(zuo)的次数。
分(fen)块处理(li)则是将输入序列分(fen)成(cheng)若(ruo)干小(xiao)块，每次只处理(li)一个小(xiao)块的数据。这种方法使(shi)得每次处理(li)的数据量减少，从(cong)而降低了内存使(shi)用和计算复杂度。
这样一来，两个关键问(wen)题(ti)就得到了解决(jue)，这两大核心思想也在本次的FlashAttention-3中(zhong)得到了继承。
但是，第一代的FlashAttention也遗留下(xia)了并行(xing)性不够强、工作(zuo)分(fen)区划(hua)分(fen)不合理(li)，以及非矩阵乘法较多(duo)（GPU计算单(dan)元处理(li)矩阵乘法比非矩阵速度更快）的问(wen)题(ti)。
针对这一问(wen)题(ti)，第二代FlashAttention通(tong)过重写(xie)softmax，减少了重新缩放操作(zuo)、边(bian)界检查和因果屏蔽操作(zuo)的次数，使(shi)得大部分(fen)计算集(ji)中(zhong)在矩阵乘法上。
另外，FlashAttention-2引入了序列长度维度上的并行(xing)化，并针对工作(zuo)在线程块之间的分(fen)配进行(xing)了优(you)化，GPU利用效(xiao)率更高了。
可以说前两代当中(zhong)，作(zuo)者一直(zhi)坚持着充分(fen)利用硬件特点这一思路，但站在今天的视角来看，对硬件的挖掘仍然不够充分(fen)。
到了这次的FlashAttention-3，由于是直(zhi)接和英伟达官方合作(zuo)，对英伟达Hopper架构特点的理(li)解更加透彻，软(ruan)硬件之间的协同进一步增(zeng)强了。
FlashAttention-3的技(ji)术(shu)报告显示，为了充分(fen)匹配Hopper架构，团队(dui)主(zhu)要(yao)做了三方面的技(ji)术(shu)升级(ji)。
首先，Hopper架构的一个重要(yao)特点是Tensor Core的异步性，FlashAttention-3针对性地提出了一种异步方式。
具(ju)体来说，FlashAttention-3引入了一种“生产(chan)者（Producer）-消费者（Consumer）”的编程模型，将注意力的计算划(hua)分(fen)为两个角色。
- “生产(chan)者”负责将数据从(cong)HBM异步加载到片上共(gong)享(xiang)内存（SMEM）。这个过程主(zhu)要(yao)利用了Hopper GPU的张量内存加速器（TMA），可以在不阻塞CUDA核心的情况下(xia)进行(xing)数据传输。
- 消费者直(zhi)接从(cong)共(gong)享(xiang)内存读取数据，并使(shi)用Tensor Core执行(xing)矩阵乘法等计算密集(ji)型任务。由于共(gong)享(xiang)内存的访问(wen)延迟远低于全局内存，消费者可以快速获取所需(xu)数据，提升计算效(xiao)率。
为了实现角色的划(hua)分(fen)，作(zuo)者引入了warp专门化技(ji)术(shu)，用不同的warp分(fen)别(bie)匹配生产(chan)者和消费者，让两者可以并行(xing)执行(xing)。
这其中(zhong)利用了Hopper架构的动态warp寄(ji)存器分(fen)配特性，通(tong)过setmaxnreg指令优(you)化了寄(ji)存器资源的利用。
为了进一步提高GPU的利用率，作(zuo)者又(you)提出了一种“乒乓调度”策略，让一个warp组执行(xing)矩阵乘法时，另一个warp组执行(xing)softmax，从(cong)而实现计算的重叠。
具(ju)体讲，FlashAttention-3使(shi)用CUDA的同步原语控制不同warp组之间的执行(xing)顺序，让不同warp组分(fen)别(bie)执行(xing)两种运算，然后像乒乓球一样交替(ti)运行(xing)。
第二大技(ji)术(shu)特点，是warp组内部GEMMs和softmax的重叠，核心奥义是重新安排计算的执行(xing)顺序以提高GPU利用率。
与乒乓调度不同，这里的计算重排处理(li)的是warp组内部的重叠，而乒乓调度更关注组间协调。
实现方式上，FlashAttention-3提出了一种两阶段GEMM-softmax流水线方案，以打破不同操作(zuo)之间的数据依赖。
- 第一阶段，当前迭代（iteration）的softmax操作(zuo)与下(xia)一个迭代的Q·K^T矩阵乘法重叠执行(xing)。
- 第二阶段，当前迭代的P·V矩阵乘法与下(xia)一个迭代的softmax操作(zuo)重叠执行(xing)。
通(tong)过引入额外的寄(ji)存器和共(gong)享(xiang)内存缓冲区，FlashAttention-3实现了跨迭代的数据传递和重用。
在每个迭代中(zhong)，Q·K^T的结果首先存储在名为S_cur的缓冲区中(zhong)，用于当前迭代的softmax计算，同时异步执行(xing)下(xia)一个迭代的Q·K^T矩阵乘法，结果存储在名为S_next的缓冲区中(zhong)。
在执行(xing)当前迭代的P·V矩阵乘法时,异步执行(xing)下(xia)一个迭代的softmax操作(zuo)，并更新S_cur和S_next缓冲区。
第三项更新，是用更低的FP8精(jing)度替(ti)代FP16。
实际上，降低数值精(jing)度是一种常见的优(you)化策略，可以显著提高GPU的计算吞吐量和能效(xiao)，Hopper GPU也引入了FP8精(jing)度的Tensor Core支持。
但是，直(zhi)接将注意力计算从(cong)FP16转换为FP8可能会引入较大的精(jing)度损失。
另外，FP8 Tensor Core对输入数据的布局也有特定(ding)的要(yao)求（K维度连(lian)续），不幸的是，注意力计算中(zhong)的输入数据存储格式（头维度连(lian)续）并不符合这样的要(yao)求。
所以FlashAttention-3首先引入了一系列内存布局转换技(ji)术(shu)，动态转置V矩阵的块，改变其连(lian)续方式，从(cong)而适配FP8 Tensor Core的布局要(yao)求。
在此基础(chu)之上，为了获得更高的计算精(jing)度，FlashAttention-3又(you)采用了分(fen)块量化和非相干处理(li)技(ji)术(shu)。
传统的量化方法通(tong)常对整个矩阵使(shi)用一个统一的缩放因子（per-tensor quantization），无法很(hen)好地适应不同区域的数值范围。
FlashAttention-3则采用了分(fen)块量化（block-wise quantization）的策略，为每个块单(dan)独设置缩放因子，更好地捕捉局部的数值分(fen)布。
非相干处理(li)（incoherent processing）技(ji)术(shu)则是通(tong)过随机正交矩阵对输入数据进行(xing)旋(xuan)转，破坏不同块之间的相干性，减少量化误差的传播。
这两项技(ji)术(shu)的结合使(shi)得FlashAttention-3在FP8精(jing)度下(xia)取得了更高的计算精(jing)度，显著优(you)于传统的量化方法。
结果，与基于传统量化方法的FP8实现相比，FlashAttention-3的使(shi)得精(jing)度提高了2.6倍。
比标准Attention快16倍
以上就是FlashAttention-3在充分(fen)研(yan)究Hopper架构特点后做出的三大更新，针对更新后的表(biao)现，作(zuo)者主(zhu)要(yao)进行(xing)了3方面测试(shi)。
- 注意力基准测试(shi)
- 消融实验
- FP8注意力准确性测试(shi)
首先来看注意力基准测试(shi)。
通(tong)过改变序列长度（512、1k、……16k），并设置批大小(xiao)以确保总(zong)token数为16k。研(yan)究人员将隐藏(cang)维度设置为2048，头维度设置为64、128或258，计算前向传播、后向传播。
对比标准Attention、FlashAttention-2、Triton、cuDNN和FlashAttention-3，在H100 80GB SXM5上FP16的运行(xing)时间。
FlashAttention-3的前向传播比FlashAttention-2快1.5-2倍，后向传播快1.5-1.75倍。
与标准Attention相比，FlashAttention-3的速度快了3-16倍。
对于中(zhong)长序列（1k以上），FlashAttention-3甚至(zhi)超过了专门为H100优(you)化的cuDNN。
在消融实验中(zhong)，通(tong)过对非因果FP16 FlashAttention-3进行(xing)了2阶段WGMMA-softmax流水线和warp特殊化的消融研(yan)究，参数固定(ding)为{batch, seqlen, nheads, hdim} = {4, 8448, 16, 128}。
结果证实，FlashAttention-3改进带来了显著加速，从(cong)570提升到661。
另外，因为对FlashAttention的数值误差感兴(xing)趣，研(yan)究团队(dui)还(hai)将FlashAttention-2、FlashAttention-3和标准Attention进行(xing)了比较。
为了模拟LLMs中(zhong)的异常特征和激活，研(yan)究团队(dui)生成(cheng)了Q、K、V的条目，分(fen)布为：N(0,1)+N(0,100)·Bernoulli(0.001)
也就是说，每个条目都服从(cong)均值为0、标准差为1的正态分(fen)布，但对于0.1%的条目，增(zeng)加了一个独立的项，其标准差为10。然后测量均方根误差（RMSE）。
结果显示，在FP16中(zhong)，由于中(zhong)间结果（softmax）保留在FP32中(zhong)，FlashAttention-2和FlashAttention-3的RMSE比标准Attention减少1.7倍。
FP8的标准Attention使(shi)用每个张量的缩放，matmul累加器在FP32中(zhong)，中(zhong)间softmax结果保留在FP16中(zhong)。由于块量化和非相干处理(li)，FP8中(zhong)的FlashAttention-3比这个基线更准确2.6倍。
最后，论文还(hai)表(biao)示目前工作(zuo)专注于Hopper架构，后续将推广到其他硬件。
除了英伟达为研(yan)究提供了技(ji)术(shu)支持外，Meta、Together AI和普林斯顿大学为研(yan)究提供了计算支持。
本文来源：量子位，原文标题(ti)：《H100利用率飙升至(zhi)75%！英伟达亲自下(xia)场FlashAttention三代升级(ji)，比标准注意力快16倍》
*免责声明：文章内容仅供参考，不构成(cheng)投资建议
*风险提示：股市有风险，入市需(xu)谨慎