亚洲综合av网_中文字幕一区二区三区在线观看_91免费看片在线观看_久久精品久久精品亚洲人

行業(yè)門戶
NEWS CENTRE
首頁
>
新聞中心
>
H100利用率飆升至75%!英偉達(dá)親自下場(chǎng)FlashAttention三代升級(jí),比標(biāo)準(zhǔn)注意力快16倍
H100利用率飆升至75%!英偉達(dá)親自下場(chǎng)FlashAttention三代升級(jí),比標(biāo)準(zhǔn)注意力快16倍
2024-07-13 閱讀:744

來源:量子位

時(shí)隔一年,F(xiàn)lashAttention-3已經(jīng)全方位升級(jí)。訓(xùn)練速度提升1.5-2倍,F(xiàn)P16下計(jì)算吞吐量高達(dá)740TFLOPs/s,達(dá)理論最大吞吐量75%,更充分利用計(jì)算資源,此前只能做到35%。FP8下速度接近1.2PFLOPs/s!同時(shí)誤差也進(jìn)一步減小,F(xiàn)P8下的誤差比標(biāo)準(zhǔn)Attention減少2.6倍。

大模型訓(xùn)練推理神作,又更新了!

主流大模型都在用的FlashAttention,剛剛升級(jí)第三代。

時(shí)隔一年,F(xiàn)lashAttention-3已經(jīng)全方位升級(jí)。

訓(xùn)練速度提升1.5-2倍,F(xiàn)P16下計(jì)算吞吐量高達(dá)740TFLOPs/s,達(dá)理論最大吞吐量75%,更充分利用計(jì)算資源,此前只能做到35%。

FP8下速度接近1.2PFLOPs/s!

同時(shí)誤差也進(jìn)一步減小,F(xiàn)P8下的誤差比標(biāo)準(zhǔn)Attention減少2.6倍

而且這一次,不再是一作Tri Dao單打獨(dú)斗,F(xiàn)lashAttention-3直接和英偉達(dá)、Meta、谷歌等合作,針對(duì)最強(qiáng)芯片H100專門做優(yōu)化。

英偉達(dá)CUTLASS團(tuán)隊(duì)和cuDNN團(tuán)隊(duì),都直接為該研究提供支持。

同時(shí)和前作一樣,F(xiàn)lashAttention-3也將開源,PyTorch和Hugging Face中都集成。

作者之一Vijay Thakkar激動(dòng)表示:

曾經(jīng)在FA2發(fā)布時(shí),我就說過這句話。今天,我想再說一次:

看到CUTLASS和CuTe被用來開讓Tensor Core大顯身手的新算法,真的泰褲辣。

前Stable Diffusion老板Emad也非常關(guān)注這一進(jìn)展,他推測(cè)使用FlashAttention-3,能將4090的FP8計(jì)算吞吐量推升到700+TFLOPs。

充分利用Hopper架構(gòu)特點(diǎn)

自初代發(fā)布以來,F(xiàn)lashAttention已經(jīng)使大模型速度提高了4-8倍,但還有一個(gè)遺憾:尚未充分利用現(xiàn)代 GPU。

針對(duì)英偉達(dá)H100倍后的Hopper架構(gòu)新特性,三代進(jìn)行了專門優(yōu)化。

整個(gè)系列的核心思路,是IO感知優(yōu)化分塊處理

作者認(rèn)為,傳統(tǒng)的注意力機(jī)制效率低的原因,在處理長(zhǎng)序列時(shí),會(huì)出現(xiàn)內(nèi)存訪問操作頻繁,以及算法復(fù)雜度指數(shù)級(jí)暴增這兩大問題。

FlashAttention通過IO感知優(yōu)化將數(shù)據(jù)從較大但緩慢的高帶寬內(nèi)存(HBM)加載到較小但更快的片上內(nèi)存(SRAM),在SRAM中執(zhí)行計(jì)算,減少了內(nèi)存讀寫操作的次數(shù)。

分塊處理則是將輸入序列分成若干小塊,每次只處理一個(gè)小塊的數(shù)據(jù)。這種方法使得每次處理的數(shù)據(jù)量減少,從而降低了內(nèi)存使用和計(jì)算復(fù)雜度。

這樣一來,兩個(gè)關(guān)鍵問題就得到了解決,這兩大核心思想也在本次的FlashAttention-3中得到了繼承。

但是,第一代的FlashAttention也遺留下了并行性不夠強(qiáng)、工作分區(qū)劃分不合理,以及非矩陣乘法較多(GPU計(jì)算單元處理矩陣乘法比非矩陣速度更快)的問題。

針對(duì)這一問題,第二代FlashAttention通過重寫softmax,減少了重新縮放操作、邊界檢查和因果屏蔽操作的次數(shù),使得大部分計(jì)算集中在矩陣乘法上。

另外,F(xiàn)lashAttention-2引入了序列長(zhǎng)度維度上的并行化,并針對(duì)工作在線程塊之間的分配進(jìn)行了優(yōu)化,GPU利用效率更高了。

可以說前兩代當(dāng)中,作者一直堅(jiān)持著充分利用硬件特點(diǎn)這一思路,但站在今天的視角來看,對(duì)硬件的挖掘仍然不夠充分。

到了這次的FlashAttention-3,由于是直接和英偉達(dá)官方合作,對(duì)英偉達(dá)Hopper架構(gòu)特點(diǎn)的理解更加透徹,軟硬件之間的協(xié)同進(jìn)一步增強(qiáng)了。

FlashAttention-3的技術(shù)報(bào)告顯示,為了充分匹配Hopper架構(gòu),團(tuán)隊(duì)主要做了三方面的技術(shù)升級(jí)。

首先,Hopper架構(gòu)的一個(gè)重要特點(diǎn)是Tensor Core的異步性,F(xiàn)lashAttention-3針對(duì)性地提出了一種異步方式。

具體來說,F(xiàn)lashAttention-3引入了一種“生產(chǎn)者(Producer)-消費(fèi)者(Consumer)”的編程模型,將注意力的計(jì)算劃分為兩個(gè)角色。

  • “生產(chǎn)者”負(fù)責(zé)將數(shù)據(jù)從HBM異步加載到片上共享內(nèi)存(SMEM)。這個(gè)過程主要利用了Hopper GPU的張量?jī)?nèi)存加速器(TMA),可以在不阻塞CUDA核心的情況下進(jìn)行數(shù)據(jù)傳輸。

  • 消費(fèi)者直接從共享內(nèi)存讀取數(shù)據(jù),并使用Tensor Core執(zhí)行矩陣乘法等計(jì)算密集型任務(wù)。由于共享內(nèi)存的訪問延遲遠(yuǎn)低于全局內(nèi)存,消費(fèi)者可以快速獲取所需數(shù)據(jù),提升計(jì)算效率。

為了實(shí)現(xiàn)角色的劃分,作者引入了warp專門化技術(shù),用不同的warp分別匹配生產(chǎn)者和消費(fèi)者,讓兩者可以并行執(zhí)行。

這其中利用了Hopper架構(gòu)的動(dòng)態(tài)warp寄存器分配特性,通過setmaxnreg指令優(yōu)化了寄存器資源的利用。

為了進(jìn)一步提高GPU的利用率,作者又提出了一種“乒乓調(diào)度”策略,讓一個(gè)warp組執(zhí)行矩陣乘法時(shí),另一個(gè)warp組執(zhí)行softmax,從而實(shí)現(xiàn)計(jì)算的重疊。

具體講,F(xiàn)lashAttention-3使用CUDA的同步原語控制不同warp組之間的執(zhí)行順序,讓不同warp組分別執(zhí)行兩種運(yùn)算,然后像乒乓球一樣交替運(yùn)行。

第二大技術(shù)特點(diǎn),是warp組內(nèi)部GEMMs和softmax的重疊,核心奧義是重新安排計(jì)算的執(zhí)行順序以提高GPU利用率。

與乒乓調(diào)度不同,這里的計(jì)算重排處理的是warp組內(nèi)部的重疊,而乒乓調(diào)度更關(guān)注組間協(xié)調(diào)。

實(shí)現(xiàn)方式上,F(xiàn)lashAttention-3提出了一種兩階段GEMM-softmax流水線方案,以打破不同操作之間的數(shù)據(jù)依賴。

  • 第一階段,當(dāng)前迭代(iteration)的softmax操作與下一個(gè)迭代的Q·K^T矩陣乘法重疊執(zhí)行。

  • 第二階段,當(dāng)前迭代的P·V矩陣乘法與下一個(gè)迭代的softmax操作重疊執(zhí)行。

通過引入額外的寄存器和共享內(nèi)存緩沖區(qū),F(xiàn)lashAttention-3實(shí)現(xiàn)了跨迭代的數(shù)據(jù)傳遞和重用。

在每個(gè)迭代中,Q·K^T的結(jié)果首先存儲(chǔ)在名為S_cur的緩沖區(qū)中,用于當(dāng)前迭代的softmax計(jì)算,同時(shí)異步執(zhí)行下一個(gè)迭代的Q·K^T矩陣乘法,結(jié)果存儲(chǔ)在名為S_next的緩沖區(qū)中。

在執(zhí)行當(dāng)前迭代的P·V矩陣乘法時(shí),異步執(zhí)行下一個(gè)迭代的softmax操作,并更新S_cur和S_next緩沖區(qū)。

第三項(xiàng)更新,是用更低的FP8精度替代FP16。

實(shí)際上,降低數(shù)值精度是一種常見的優(yōu)化策略,可以顯著提高GPU的計(jì)算吞吐量和能效,Hopper GPU也引入了FP8精度的Tensor Core支持。

但是,直接將注意力計(jì)算從FP16轉(zhuǎn)換為FP8可能會(huì)引入較大的精度損失。

另外,F(xiàn)P8 Tensor Core對(duì)輸入數(shù)據(jù)的布局也有特定的要求(K維度連續(xù)),不幸的是,注意力計(jì)算中的輸入數(shù)據(jù)存儲(chǔ)格式(頭維度連續(xù))并不符合這樣的要求。

所以FlashAttention-3首先引入了一系列內(nèi)存布局轉(zhuǎn)換技術(shù),動(dòng)態(tài)轉(zhuǎn)置V矩陣的塊,改變其連續(xù)方式,從而適配FP8 Tensor Core的布局要求。

在此基礎(chǔ)之上,為了獲得更高的計(jì)算精度,F(xiàn)lashAttention-3又采用了分塊量化非相干處理技術(shù)。

傳統(tǒng)的量化方法通常對(duì)整個(gè)矩陣使用一個(gè)統(tǒng)一的縮放因子(per-tensor quantization),無法很好地適應(yīng)不同區(qū)域的數(shù)值范圍。

FlashAttention-3則采用了分塊量化(block-wise quantization)的策略,為每個(gè)塊單獨(dú)設(shè)置縮放因子,更好地捕捉局部的數(shù)值分布。

非相干處理(incoherent processing)技術(shù)則是通過隨機(jī)正交矩陣對(duì)輸入數(shù)據(jù)進(jìn)行旋轉(zhuǎn),破壞不同塊之間的相干性,減少量化誤差的傳播。

這兩項(xiàng)技術(shù)的結(jié)合使得FlashAttention-3在FP8精度下取得了更高的計(jì)算精度,顯著優(yōu)于傳統(tǒng)的量化方法。

結(jié)果,與基于傳統(tǒng)量化方法的FP8實(shí)現(xiàn)相比,F(xiàn)lashAttention-3的使得精度提高了2.6倍。

比標(biāo)準(zhǔn)Attention快16倍

以上就是FlashAttention-3在充分研究Hopper架構(gòu)特點(diǎn)后做出的三大更新,針對(duì)更新后的表現(xiàn),作者主要進(jìn)行了3方面測(cè)試。

  • 注意力基準(zhǔn)測(cè)試

  • 消融實(shí)驗(yàn)

  • FP8注意力準(zhǔn)確性測(cè)試

首先來看注意力基準(zhǔn)測(cè)試。

通過改變序列長(zhǎng)度(512、1k、……16k),并設(shè)置批大小以確保總token數(shù)為16k。研究人員將隱藏維度設(shè)置為2048,頭維度設(shè)置為64、128或258,計(jì)算前向傳播、后向傳播。

對(duì)比標(biāo)準(zhǔn)Attention、FlashAttention-2、Triton、cuDNN和FlashAttention-3,在H100 80GB SXM5上FP16的運(yùn)行時(shí)間。

FlashAttention-3的前向傳播比FlashAttention-2快1.5-2倍,后向傳播快1.5-1.75倍。

與標(biāo)準(zhǔn)Attention相比,F(xiàn)lashAttention-3的速度快了3-16倍。

對(duì)于中長(zhǎng)序列(1k以上),F(xiàn)lashAttention-3甚至超過了專門為H100優(yōu)化的cuDNN。

在消融實(shí)驗(yàn)中,通過對(duì)非因果FP16 FlashAttention-3進(jìn)行了2階段WGMMA-softmax流水線和warp特殊化的消融研究,參數(shù)固定為{batch, seqlen, nheads, hdim} = {4, 8448, 16, 128}。

結(jié)果證實(shí),F(xiàn)lashAttention-3改進(jìn)帶來了顯著加速,從570提升到661。

另外,因?yàn)閷?duì)FlashAttention的數(shù)值誤差感興趣,研究團(tuán)隊(duì)還將FlashAttention-2、FlashAttention-3和標(biāo)準(zhǔn)Attention進(jìn)行了比較。

為了模擬LLMs中的異常特征和激活,研究團(tuán)隊(duì)生成了Q、K、V的條目,分布為:N(0,1)+N(0,100)?Bernoulli(0.001)

也就是說,每個(gè)條目都服從均值為0、標(biāo)準(zhǔn)差為1的正態(tài)分布,但對(duì)于0.1%的條目,增加了一個(gè)獨(dú)立的項(xiàng),其標(biāo)準(zhǔn)差為10。然后測(cè)量均方根誤差(RMSE)。

結(jié)果顯示,在FP16中,由于中間結(jié)果(softmax)保留在FP32中,F(xiàn)lashAttention-2和FlashAttention-3的RMSE比標(biāo)準(zhǔn)Attention減少1.7倍

FP8的標(biāo)準(zhǔn)Attention使用每個(gè)張量的縮放,matmul累加器在FP32中,中間softmax結(jié)果保留在FP16中。由于塊量化和非相干處理,F(xiàn)P8中的FlashAttention-3比這個(gè)基線更準(zhǔn)確2.6倍

最后,論文還表示目前工作專注于Hopper架構(gòu),后續(xù)將推廣到其他硬件。

除了英偉達(dá)為研究提供了技術(shù)支持外,Meta、Together AI和普林斯頓大學(xué)為研究提供了計(jì)算支持。

本文來源:量子位,原文標(biāo)題:《H100利用率飆升至75%!英偉達(dá)親自下場(chǎng)FlashAttention三代升級(jí),比標(biāo)準(zhǔn)注意力快16倍》

風(fēng)險(xiǎn)提示及免責(zé)條款

市場(chǎng)有風(fēng)險(xiǎn),投資需謹(jǐn)慎。本文不構(gòu)成個(gè)人投資建議,也未考慮到個(gè)別用戶特殊的投資目標(biāo)、財(cái)務(wù)狀況或需要。用戶應(yīng)考慮本文中的任何意見、觀點(diǎn)或結(jié)論是否符合其特定狀況。據(jù)此投資,責(zé)任自負(fù)。

13560189272
地址:廣州市天河區(qū)黃埔大道西201號(hào)金澤大廈808室
COPYRIFHT ? 2010-2020 廣州市名聯(lián)網(wǎng)絡(luò)科技有限公司 ALL RIGHTS RESERVED 粵ICP備10203057號(hào)
  • 這里是二維碼
亚洲综合av网_中文字幕一区二区三区在线观看_91免费看片在线观看_久久精品久久精品亚洲人
蜜桃av一区二区| 91黄色免费看| 亚洲激情在线播放| 中文字幕久久午夜不卡| 久久久精品免费观看| 精品美女一区二区三区| 日韩精品中文字幕在线一区| 制服丝袜激情欧洲亚洲| 7777精品伊人久久久大香线蕉最新版 | 午夜欧美视频在线观看| 洋洋成人永久网站入口| 亚洲国产成人av好男人在线观看| 亚洲福利一二三区| 日本不卡123| 国产一区在线视频| 成人av集中营| 欧美性猛交xxxxxx富婆| 欧美一级xxx| 国产偷v国产偷v亚洲高清| 国产精品久久久久久妇女6080 | 综合网在线视频| 亚洲精品乱码久久久久久久久| 一级中文字幕一区二区| 婷婷开心激情综合| 国内精品免费在线观看| av在线综合网| 制服丝袜av成人在线看| 国产清纯白嫩初高生在线观看91 | 激情五月婷婷综合| caoporen国产精品视频| 欧美午夜寂寞影院| 日韩欧美国产一区二区三区| 欧美国产精品一区二区| 一二三区精品福利视频| 蜜桃精品在线观看| 91色视频在线| 精品国产污网站| 一区二区三区视频在线看| 麻豆精品一二三| 91免费视频观看| 精品乱人伦一区二区三区| 亚洲老司机在线| 国产另类ts人妖一区二区| 欧洲在线/亚洲| 国产欧美综合在线观看第十页| 亚洲第一狼人社区| www.亚洲精品| 欧美xxx久久| 亚洲大片一区二区三区| 丰满白嫩尤物一区二区| 欧美一激情一区二区三区| 亚洲女同ⅹxx女同tv| 国产传媒欧美日韩成人| 日韩视频一区二区在线观看| 亚洲综合色婷婷| 成人sese在线| 久久精品欧美日韩| 久久精品国产成人一区二区三区 | 国产精品一区专区| 日韩欧美国产精品| 亚洲一区二区欧美激情| 北条麻妃国产九九精品视频| 久久久久久久性| 日本人妖一区二区| 欧美在线不卡一区| 中文字幕色av一区二区三区| 国产一区二区女| 欧美va亚洲va在线观看蝴蝶网| 亚洲午夜电影在线| 色综合 综合色| 日韩毛片一二三区| 成人深夜在线观看| 久久久久久久久久久久久夜| 蜜桃一区二区三区在线观看| 4hu四虎永久在线影院成人| 亚洲综合在线电影| 色哟哟一区二区在线观看| 国产精品免费人成网站| 国产成人aaaa| 国产精品美女久久久久久| 国产福利一区在线| 国产精品女人毛片| av不卡免费电影| 亚洲视频资源在线| 欧美午夜一区二区| 天天色天天爱天天射综合| 欧美日韩国产免费| 日本一区中文字幕 | 一区二区在线免费观看| 91福利视频在线| 国产精品一区二区91| 视频一区二区不卡| 99久久久国产精品免费蜜臀| 国产精品久久久久久久久搜平片 | 奇米影视一区二区三区小说| 日韩欧美一区在线| 九色|91porny| 久久综合色8888| 风间由美一区二区三区在线观看| 国产精品国产自产拍高清av | 久久精品视频网| 成人一区二区三区视频| 国产精品久久久久久一区二区三区| 99精品桃花视频在线观看| 一区二区三区欧美亚洲| 在线成人免费观看| 国产在线播放一区| 亚洲欧美日韩国产综合在线 | 成人免费一区二区三区在线观看| 91亚洲男人天堂| 亚洲成a人片在线不卡一二三区| 日韩免费观看2025年上映的电影 | 日韩精品一区二| av不卡免费电影| 日精品一区二区| 久久久久9999亚洲精品| 欧美亚洲一区三区| 韩国女主播成人在线| 亚洲理论在线观看| 久久免费偷拍视频| 色94色欧美sute亚洲线路二| 久久精品国产**网站演员| 中文字幕在线不卡国产视频| 91精品国产欧美日韩| 94-欧美-setu| 久久99最新地址| 一区二区三区在线影院| 26uuu成人网一区二区三区| 91黄色免费网站| 国产91露脸合集magnet| 青青草原综合久久大伊人精品| 国产精品福利在线播放| 91精品国产一区二区人妖| 91香蕉国产在线观看软件| 韩国v欧美v日本v亚洲v| 亚洲va国产天堂va久久en| 国产精品国产三级国产普通话99| 91精品欧美综合在线观看最新| www.色综合.com| 国产一区激情在线| 麻豆极品一区二区三区| 亚洲国产精品一区二区尤物区| 国产人妖乱国产精品人妖| 337p亚洲精品色噜噜噜| 91高清在线观看| 91在线看国产| 成人免费福利片| 国产精品亚洲人在线观看| 免费人成网站在线观看欧美高清| 亚洲综合一二三区| 亚洲欧美日韩人成在线播放| 国产精品毛片大码女人| 国产视频在线观看一区二区三区| 91精选在线观看| 欧美色精品在线视频| 91视频在线观看| 不卡av免费在线观看| 国产成人啪免费观看软件 | 欧美日韩国产乱码电影| 欧美性高清videossexo| 色诱亚洲精品久久久久久| 成人国产精品免费观看| 成人三级在线视频| 成人午夜免费电影| 成人午夜又粗又硬又大| 国产99一区视频免费| 国产精品88av| 成人激情开心网| 99国产精品国产精品毛片| 91最新地址在线播放| 色欧美88888久久久久久影院| 日本精品裸体写真集在线观看 | 免费在线观看日韩欧美| 免费不卡在线视频| 看电影不卡的网站| 韩国一区二区在线观看| 国产高清不卡一区二区| 成人免费视频免费观看| 北条麻妃国产九九精品视频| 色综合久久天天| 欧美电影一区二区三区| 欧美大胆人体bbbb| 国产精品午夜久久| 亚洲免费av网站| 日韩国产欧美一区二区三区| 久久超碰97中文字幕| 风间由美性色一区二区三区| 91美女蜜桃在线| 69p69国产精品| 久久人人97超碰com| ㊣最新国产の精品bt伙计久久| 亚洲一区二区黄色| 久久成人精品无人区| 播五月开心婷婷综合| 欧美日韩精品电影| 久久先锋影音av| 一片黄亚洲嫩模| 国产真实乱偷精品视频免| 91小视频免费看| 精品久久久久久久久久久院品网|