化實(shí)戰(zhàn)：從Tiling到端側(cè)部署的全鏈路加速)
1. 為什么算子優(yōu)化才是全鏈路加速的命門先聊一個(gè)很多剛接觸異構(gòu)計(jì)算的同行容易忽略的事實(shí)一套加速鏈路里最耗時(shí)間的往往不是模型結(jié)構(gòu)本身而是算子層執(zhí)行效率。模型結(jié)構(gòu)是算法設(shè)計(jì)者定好的骨架算子是把骨架變成真實(shí)算力的肌肉。肌肉乏力骨架再合理也跑不出性能。我最早接觸CANNCompute Architecture for Neural Networks異構(gòu)計(jì)算架構(gòu)時(shí)也踩過這個(gè)認(rèn)知偏差拿到一個(gè)模型第一步就往圖優(yōu)化和部署工具上使勁覺得算子實(shí)現(xiàn)是芯片廠商的事自己寫好算法就行。后來在一次真實(shí)項(xiàng)目里一個(gè)簡單的自定義算子因?yàn)闆]有做Tiling切分端側(cè)推理整條鏈路延遲多了將近一倍才明白算子層是整個(gè)加速流水線的底座。底座不牢后面所有優(yōu)化都是錦上添花成不了雪中送炭。這篇文章圍繞一條完整鏈路來寫從算子的計(jì)算分析、Tiling設(shè)計(jì)、指令選型到圖編譯階段的算子融合策略再到端側(cè)部署時(shí)的模型轉(zhuǎn)換與推理配置最后落到性能調(diào)優(yōu)的閉環(huán)方法。適合已經(jīng)能跑通基礎(chǔ)推理流程、想進(jìn)一步壓榨設(shè)備算力的開發(fā)者也適合剛接觸異構(gòu)編程、想建立全局性能觀的新手。整篇文章不是照搬官方文檔而是把我實(shí)際跑過、調(diào)過、踩過坑的經(jīng)驗(yàn)拆開來講。先明確一個(gè)核心認(rèn)知全鏈路加速不是某一環(huán)的孤軍奮戰(zhàn)而是計(jì)算密度、數(shù)據(jù)搬運(yùn)、內(nèi)存布局、調(diào)度策略四個(gè)維度的協(xié)同。這四件事在算子開發(fā)階段就要想清楚拖到部署階段再補(bǔ)救空間就很有限了。2. 從計(jì)算特征出發(fā)先搞清楚你這算子到底缺什么2.1 算子性能瓶頸的判斷方法拿計(jì)算量和訪存量說話不管下游用的是什么加速芯片算子的性能瓶頸就兩種算力吃滿但數(shù)據(jù)喂不飽叫內(nèi)存瓶頸數(shù)據(jù)搬運(yùn)沒問題但計(jì)算單元在空轉(zhuǎn)叫計(jì)算瓶頸。判斷方法不復(fù)雜先算兩個(gè)數(shù)。以矩陣乘法為例一個(gè) ( M \times K ) 的矩陣乘以 ( K \times N ) 的矩陣計(jì)算量是 ( 2 \times M \times N \times K ) 次浮點(diǎn)操作數(shù)據(jù)搬運(yùn)量是 ( (M \times K K \times N M \times N) \times 4 ) 字節(jié)按FP32算。計(jì)算密度 總計(jì)算量 / 總數(shù)據(jù)量這個(gè)值如果遠(yuǎn)低于設(shè)備的計(jì)算訪存比就是內(nèi)存瓶頸如果高出很多就優(yōu)先優(yōu)化計(jì)算指令的流水線效率。我自己習(xí)慣在動(dòng)手寫算子前先做這個(gè)估算再對(duì)照設(shè)備的理論峰值來判斷優(yōu)化方向。這一步不值錢但能讓優(yōu)化避開南轅北轍的坑。有一個(gè)實(shí)際例子某模型里的歸一化算子輸入是 ( 2048 \times 2048 ) 的矩陣逐元素操作計(jì)算量很小數(shù)據(jù)量很大典型的內(nèi)存瓶頸。這種情況下拼命優(yōu)化計(jì)算指令根本沒用關(guān)鍵是減少重復(fù)搬運(yùn)讓數(shù)據(jù)在片上多復(fù)用。2.2 算子的計(jì)算模式分類決定了它適合走哪條路異構(gòu)芯片雖然指令集各有不同但算子按計(jì)算模式大致能分成幾類分類的意義在于決定優(yōu)化策略的起點(diǎn)逐元素類算子數(shù)據(jù)每個(gè)點(diǎn)獨(dú)立運(yùn)算比如激活函數(shù)、歸一化、逐元素加減。這類算子好寫但極易被內(nèi)存帶寬卡死優(yōu)化重點(diǎn)在數(shù)據(jù)搬運(yùn)和并行切分。歸約類算子需要跨行或跨列聚合比如求和、求均值、Softmax。這類算子的問題在于數(shù)據(jù)依賴需要處理好片上累加和分段歸約。卷積/矩陣乘類算子計(jì)算密度高數(shù)據(jù)復(fù)用性強(qiáng)是最容易榨出峰值算力的類型但Tiling和搬運(yùn)策略也最復(fù)雜。訪存密集類算子比如Transpose、Reshape本身計(jì)算量極小純粹搬運(yùn)優(yōu)化重點(diǎn)是讓搬運(yùn)模式連續(xù)化避免小碎片傳輸。我遇到很多開發(fā)者拿到一個(gè)陌生算子不看計(jì)算特征就先套模板寫寫完了再回頭調(diào)性能效率非常低。正確順序是反過來的先分類再定策略最后才動(dòng)手寫內(nèi)核代碼。確定策略的過程聽起來抽象其實(shí)真正操作時(shí)只需要回答三個(gè)問題數(shù)據(jù)怎么切每一塊放哪一級(jí)緩存計(jì)算用哪種指令完成這三個(gè)問題就是整個(gè)算子開發(fā)階段的核心也是接下來要展開的Tiling設(shè)計(jì)。3. Tiling設(shè)計(jì)把大矩陣切成剛好能塞進(jìn)片上內(nèi)存的形狀3.1 為什么Tiling是算子性能的第一道分水嶺異構(gòu)芯片的存儲(chǔ)層級(jí)和CPU不太一樣通常有全局內(nèi)存相當(dāng)于系統(tǒng)內(nèi)存、片上緩存相當(dāng)于L1/L2、以及緊挨計(jì)算單元的緩沖區(qū)。片上內(nèi)存容量很有限大矩陣不可能一次全放進(jìn)去必須切塊搬運(yùn)算完一塊再搬下一塊。這個(gè)切分策略就是Tiling。Tiling切得好不好直接決定兩層?xùn)|西一是數(shù)據(jù)搬運(yùn)次數(shù)切得太大放不下切得太小搬運(yùn)次數(shù)爆炸二是計(jì)算單元利用率每塊數(shù)據(jù)計(jì)算時(shí)間必須遠(yuǎn)大于搬運(yùn)時(shí)間否則計(jì)算單元大量空轉(zhuǎn)。Tiling里的核心參數(shù)是Tiling Size——在片上內(nèi)存允許的容量內(nèi)盡量讓每一塊算得久一點(diǎn)。比如片上緩沖區(qū)是256KBFP32占4字節(jié)理論上能放65536個(gè)數(shù)。如果做 ( 128 \times 128 ) 的矩陣乘一塊剛好能放下。有兩個(gè)實(shí)際的切分策略固定分塊和自適應(yīng)分塊。固定分塊適合形狀穩(wěn)定的模型代碼簡單自適應(yīng)分塊會(huì)根據(jù)輸入尺寸動(dòng)態(tài)算切法適合輸入形狀變化頻繁的場(chǎng)景。3.2 一個(gè)具體案例512×512矩陣乘的Tiling推導(dǎo)過程光說理論容易飄拿實(shí)際例子算一遍就清楚了。假設(shè)要做 ( C A \times B )A是 ( 512 \times 512 )B也是 ( 512 \times 512 )結(jié)果C同樣是 ( 512 \times 512 )。片上緩沖區(qū)我這里指的是緊挨計(jì)算單元的那個(gè)高速緩沖區(qū)假設(shè)總?cè)萘渴?28KB也就是32768個(gè)FP32數(shù)。如果直接整塊計(jì)算A要占262144個(gè)數(shù)顯然放不下。這時(shí)候按行切分把A按行切成若干塊比如每次取 ( 128 \times 512 ) 的子塊占65536個(gè)數(shù)還是超過32768。繼續(xù)調(diào)整取 ( 64 \times 512 ) 的子塊占32768個(gè)數(shù)剛好卡在容量上限。B不動(dòng)C按行累加。這樣每次搬運(yùn)A的子塊 ( 64 \times 512 32768 ) 個(gè)數(shù)128KBB還是整塊 ( 512 \times 512 262144 ) 個(gè)數(shù)1MB。B在一次算子執(zhí)行里要被反復(fù)讀取8次如果B也太大還要做列切分。從Tiling切分維度來看行切分屬于第一維切分可能還需要第二維切分來降低中間結(jié)果的存儲(chǔ)壓力。這里的關(guān)鍵是把搬運(yùn)量和復(fù)用次數(shù)放到一起看A每次只搬一小條復(fù)用得越充分搬運(yùn)總量越小B必須整體搬屬于不可少的固定開銷。實(shí)踐里我習(xí)慣每塊數(shù)據(jù)在片上至少復(fù)用幾十次才覺得劃算如果復(fù)用次數(shù)太低往往意味著切塊太小浪費(fèi)了片上存儲(chǔ)。3.3 多核切分與負(fù)載均衡別讓一個(gè)核閑著等別人除了片上內(nèi)存維度的切分還有多核維度的切分。設(shè)備上通常有多個(gè)計(jì)算核心每個(gè)核心負(fù)責(zé)一部分?jǐn)?shù)據(jù)。這里最常踩的坑是負(fù)載不均有的核數(shù)據(jù)多、算得慢整個(gè)算子等著它有的核數(shù)據(jù)少、早早算完干瞪眼。避免負(fù)載不均的核心原則是讓每個(gè)核拿到的工作量盡量一致。對(duì)于形狀規(guī)整的矩陣乘按行均分通常就夠了對(duì)于元素稀疏的算子比如某些稀疏計(jì)算數(shù)據(jù)量差異大就需要按非零元素?cái)?shù)量來分配。多核協(xié)作時(shí)的同步設(shè)計(jì)也很關(guān)鍵核與核之間需要通信時(shí)一般走全局內(nèi)存做中轉(zhuǎn)或者用芯片提供的同步機(jī)制。這里需要特別留意不要在大循環(huán)內(nèi)部頻繁同步同步本身就是開銷次數(shù)越少越好。我做過一次實(shí)驗(yàn)同樣的算子把循環(huán)內(nèi)的同步次數(shù)從8次降到2次整體性能提升了大約15%。這種收益不需要改計(jì)算邏輯只需要調(diào)整循環(huán)結(jié)構(gòu)讓每次同步之間有更多獨(dú)立計(jì)算。4. 指令選擇與數(shù)據(jù)搬運(yùn)算得快不如算得巧4.1 計(jì)算指令和搬運(yùn)指令怎么配合才能不打空轉(zhuǎn)Tiling把數(shù)據(jù)切好之后接下來就是真正的指令執(zhí)行。CANN體系下的AI Core編程計(jì)算指令大致分為三種類型矢量指令Vector處理逐元素運(yùn)算面向向量化計(jì)算矩陣指令Matrix/Cube處理矩陣乘加操作計(jì)算密度最高標(biāo)量指令Scalar處理地址計(jì)算、循環(huán)控制等輔助邏輯。真實(shí)項(xiàng)目里寫算子時(shí)最容易忽略的是標(biāo)量指令對(duì)性能的影響。我曾經(jīng)優(yōu)化過一個(gè)算子計(jì)算邏輯本身已經(jīng)用矩陣指令寫得很高效但控制流里用了大量標(biāo)量指令做地址計(jì)算結(jié)果標(biāo)量單元成了瓶頸整體耗時(shí)反而比優(yōu)化前更差。后來把地址計(jì)算提前到搬運(yùn)階段做好緩存循環(huán)內(nèi)只做查表性能才上來。結(jié)論是主循環(huán)里盡量用計(jì)算類指令少放標(biāo)量指令。指令選擇的核心原則是能用矩陣指令解決的不要用矢量指令能用矢量指令解決的不要用標(biāo)量指令一條條算。三種指令之間不是替代關(guān)系而是流水線關(guān)系——標(biāo)量做指揮矢量做批量矩陣做重火力。寫出高效算子的人往往是讓這三者在時(shí)間上重疊數(shù)據(jù)搬運(yùn)的同時(shí)上一塊數(shù)據(jù)正在計(jì)算這一塊數(shù)據(jù)正在準(zhǔn)備搬運(yùn)。4.2 數(shù)據(jù)搬運(yùn)的隱藏成本Copy是性能刺客所有異構(gòu)計(jì)算里數(shù)據(jù)搬運(yùn)的代價(jià)都被嚴(yán)重低估。搬運(yùn)指令本身執(zhí)行不慢慢的是搬運(yùn)前后的等待和地址對(duì)齊問題。我習(xí)慣把數(shù)據(jù)搬運(yùn)分成兩種塊搬運(yùn)整塊連續(xù)數(shù)據(jù)和跳變搬運(yùn)按一定間隔取數(shù)。塊搬運(yùn)效率遠(yuǎn)高于跳變搬運(yùn)這個(gè)差距在某些硬件上能到3倍以上。所以優(yōu)化搬運(yùn)的第一準(zhǔn)則就是讓搬運(yùn)盡量連續(xù)。如果輸入數(shù)據(jù)的原始排布不連續(xù)先做一次重排再搬運(yùn)連續(xù)塊往往比直接多次跳變搬運(yùn)更快。這個(gè)先重排再搬運(yùn)的思路有點(diǎn)反直覺因?yàn)槎嗔艘淮巫x寫但實(shí)測(cè)下來很多場(chǎng)景是賺的。原因是連續(xù)搬運(yùn)能充分利用總線帶寬跳變搬運(yùn)每次都浪費(fèi)大量總線位寬。數(shù)據(jù)搬運(yùn)還有對(duì)齊要求。主流加速芯片普遍要求搬運(yùn)的起始地址和長度按特定字節(jié)對(duì)齊通常是32字節(jié)有些甚至是64字節(jié)。不對(duì)齊的搬運(yùn)要么報(bào)錯(cuò)要么被硬件強(qiáng)行拆分性能直接下滑。所以算子開發(fā)階段通常會(huì)有對(duì)齊這個(gè)參數(shù)它不是一個(gè)可配可不配的開關(guān)而是直接影響搬運(yùn)效率的關(guān)鍵條件。4.3 矢量指令寫算子一個(gè)ReLU算子的完整代碼拆解為了把上面的概念落到代碼層面我用Ascend C寫一個(gè)最簡單的ReLU算子展示矢量指令的基本用法和結(jié)構(gòu)#include kernel_operator.h using namespace AscendC; constexpr int32_t BUFFER_NUM 2; // 雙緩沖 class KernelReLU { public: __aicore__ inline KernelReLU() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, int32_t totalLen) { this-totalLen totalLen; this-blockLength totalLen / GetBlockNum(); // 每核處理長度 xGm.SetGlobalBuffer((__gm__ half*)x GetBlockIdx() * blockLength, blockLength); yGm.SetGlobalBuffer((__gm__ half*)y GetBlockIdx() * blockLength, blockLength); pipe.InitBuffer(inQueue, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(outQueue, BUFFER_NUM, blockLength * sizeof(half)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { LocalTensorhalf inLocal inQueue.AllocTensorhalf(); DataCopy(inLocal, xGm, blockLength); inQueue.EnQue(inLocal); } __aicore__ inline void Compute() { LocalTensorhalf inLocal inQueue.DeQuehalf(); LocalTensorhalf outLocal outQueue.AllocTensorhalf(); Maximum(outLocal, inLocal, half(0), blockLength); // 矢量指令取最大值等效ReLU outQueue.EnQue(outLocal); inQueue.FreeTensor(inLocal); } __aicore__ inline void CopyOut() { LocalTensorhalf outLocal outQueue.DeQuehalf(); DataCopy(yGm, outLocal, blockLength); outQueue.FreeTensor(outLocal); } private: TPipe pipe; TGlobalBufferhalf xGm, yGm; TQueQuePosition::VECIN, BUFFER_NUM inQueue; TQueQuePosition::VECOUT, BUFFER_NUM outQueue; int32_t totalLen 0; int32_t blockLength 0; }; extern C __global__ __aicore__ void ReLU(GM_ADDR x, GM_ADDR y, int32_t totalLen) { KernelReLU op; op.Init(x, y, totalLen); op.Process(); }這段代碼里已經(jīng)包含了雙緩沖結(jié)構(gòu)inQueue和outQueue各自有2個(gè)緩沖區(qū)。處理當(dāng)前數(shù)據(jù)塊的同時(shí)下一個(gè)數(shù)據(jù)塊已經(jīng)開始搬運(yùn)搬運(yùn)和計(jì)算重疊起來隱藏了大部分搬運(yùn)延遲。Maximum指令一次處理半精度浮點(diǎn)數(shù)直接走矢量單元語義上等價(jià)于ReLU因?yàn)樗龅氖虑榫褪前沿?fù)值壓到0。這個(gè)例子里有個(gè)細(xì)節(jié)值得留意緩沖區(qū)數(shù)量設(shè)成2是實(shí)現(xiàn)流水線重疊的最基本配置。如果設(shè)成1搬運(yùn)和計(jì)算就必須嚴(yán)格串行性能至少差一半。這是算子開發(fā)里雙緩沖設(shè)計(jì)的關(guān)鍵意義。5. 圖編譯與算子融合讓模型在跑之前先瘦身5.1 編譯期到底干了什么構(gòu)圖、優(yōu)化、內(nèi)存分配三步走算子單獨(dú)寫好了但模型不是一個(gè)個(gè)算子孤立跑的。CANN把整張模型圖交給編譯器統(tǒng)一編排這個(gè)階段通常叫圖編譯。圖編譯的核心工作大致可以分成三步構(gòu)圖與解析、圖優(yōu)化與融合、內(nèi)存分配與調(diào)度。構(gòu)圖階段把模型的結(jié)構(gòu)描述轉(zhuǎn)換成內(nèi)部圖表示這一步會(huì)做算子合法性檢查比如輸入輸出的數(shù)據(jù)類型是否匹配、維度是否符合算子約束。圖優(yōu)化階段是性能的關(guān)鍵包含算子融合、常量折疊、死分支消除等手段。內(nèi)存分配階段決定每個(gè)中間張量放到哪塊物理內(nèi)存、什么時(shí)候釋放做得好的話能大幅降低峰值內(nèi)存占用。我見過不少開發(fā)者繞開圖編譯階段直接把每個(gè)算子單獨(dú)執(zhí)行覺得那樣控制力更強(qiáng)。如果只有一個(gè)算子可能無所謂但在模型級(jí)別這樣做等于放棄了編譯器層面的自動(dòng)優(yōu)化——算子融合、內(nèi)存復(fù)用這些都享受不到了。正確的做法是自定義算子先單獨(dú)驗(yàn)證正確性然后放進(jìn)整個(gè)圖里一起編譯讓編譯器統(tǒng)一編排。5.2 算子融合的具體收益一個(gè)實(shí)測(cè)數(shù)據(jù)算子融合是圖優(yōu)化里收益最直觀的手段。原理不復(fù)雜把多個(gè)可合并成一個(gè)計(jì)算過程的算子融合成一個(gè)算子減少中間張量的落盤和讀回。舉一個(gè)實(shí)際案例。某視覺模型里有一段連續(xù)的卷積、批歸一化、ReLU三個(gè)操作。未融合時(shí)卷積的輸出要寫回全局內(nèi)存批歸一化再從全局內(nèi)存讀出來算完再寫回ReLU再讀一次。三次搬運(yùn)每個(gè)中間結(jié)果都是一整張?zhí)卣鲌D。融合之后卷積計(jì)算的結(jié)果直接留在片上緩沖區(qū)批歸一化和ReLU在這個(gè)結(jié)果上原位計(jì)算最后只寫一次最終結(jié)果。實(shí)測(cè)下來這個(gè)融合把三段串行算子的總耗時(shí)壓縮了約43%主要省的就是中間結(jié)果反復(fù)搬運(yùn)的時(shí)間。這里要特別提一句融合不是所有算子都能做必須滿足數(shù)據(jù)依賴關(guān)系是線性的、中間結(jié)果不需要給其他分支使用的條件。強(qiáng)行融合會(huì)破壞數(shù)據(jù)依賴或增加寄存器壓力得不償失。5.3 自動(dòng)混合精度與內(nèi)存復(fù)用編譯器的隱形加速器除了算子融合圖編譯階段還有兩個(gè)容易被忽視的優(yōu)化手段自動(dòng)混合精度和內(nèi)存復(fù)用。自動(dòng)混合精度是把一部分算子從FP32降到FP16執(zhí)行在幾乎不影響精度的情況下把計(jì)算量減半、搬運(yùn)量減半。但這不是無腦全降某些對(duì)數(shù)值敏感的操作比如Loss計(jì)算、部分歸一化層必須保留FP32。CANN的自動(dòng)混合精度機(jī)制有算子白名單和黑名單策略能自動(dòng)識(shí)別哪些算子適合降精度。內(nèi)存復(fù)用則是把生命周期不重疊的中間張量分配到同一塊物理內(nèi)存。比如兩個(gè)算子的輸出不會(huì)同時(shí)存在它們就可以共享內(nèi)存。這個(gè)優(yōu)化對(duì)端側(cè)部署尤其重要因?yàn)槎藗?cè)內(nèi)存非常有限復(fù)用得好的模型能省出20%到30%的內(nèi)存占用直接關(guān)系到模型能不能在這臺(tái)設(shè)備上跑起來。6. 端側(cè)部署從模型文件到高效推理服務(wù)的完整鏈路6.1 ATC模型轉(zhuǎn)換時(shí)那些容易被忽略的配置項(xiàng)模型在PC端訓(xùn)練好之后不能直接拿到端側(cè)設(shè)備上跑需要經(jīng)過一步模型轉(zhuǎn)換。在CANN體系里這個(gè)轉(zhuǎn)換工具通常叫ATCAscend Tensor Compiler負(fù)責(zé)把訓(xùn)練框架的模型格式轉(zhuǎn)換成端側(cè)推理設(shè)備能執(zhí)行的格式。ATC轉(zhuǎn)換的核心輸入包括原始模型文件、輸入數(shù)據(jù)的形狀和格式定義、輸出節(jié)點(diǎn)名、目標(biāo)設(shè)備型號(hào)。這幾個(gè)信息一個(gè)都不能錯(cuò)。我最常踩的坑是輸入形狀定義不匹配訓(xùn)練時(shí)模型是動(dòng)態(tài)Shape輸入尺寸每次都可能變但轉(zhuǎn)換時(shí)如果不顯式指定Shape范圍轉(zhuǎn)換出來的模型很可能只支持固定尺寸推理時(shí)一旦輸入尺寸變了直接報(bào)錯(cuò)。端側(cè)部署的常見場(chǎng)景里輸入形狀可能是固定大小也可能在一定范圍內(nèi)浮動(dòng)把握好這兩類需求后就能在模板和邊界條件之間做好取舍。如果需要在推理階段動(dòng)態(tài)改變輸入尺寸就得配置動(dòng)態(tài)Shape功能同時(shí)設(shè)置好Shape范圍。范圍設(shè)太大會(huì)浪費(fèi)內(nèi)存和算力范圍設(shè)太小超出范圍的輸入會(huì)被拒絕。合理的做法是根據(jù)真實(shí)業(yè)務(wù)場(chǎng)景統(tǒng)計(jì)輸入尺寸分布覆蓋絕大多數(shù)情況即可沒必要追求理論上的全范圍支持。轉(zhuǎn)換后生成的文件一般包含模型結(jié)構(gòu)和權(quán)重?cái)?shù)據(jù)是端側(cè)推理引擎直接加載的格式。轉(zhuǎn)換過程還會(huì)做一次圖優(yōu)化把前面提到的融合、量化等策略固化到最終模型里所以ATC階段對(duì)推理性能的影響不亞于算子開發(fā)階段。6.2 ACL推理流程的最小可跑骨架模型轉(zhuǎn)換完成后端側(cè)推理一般通過ACLAscendCL推理API來執(zhí)行。這套API的設(shè)計(jì)思路和很多推理運(yùn)行時(shí)類似初始化設(shè)備、加載模型、創(chuàng)建輸入輸出、執(zhí)行推理、釋放資源。最小可跑的推理流程骨架如下// 初始化設(shè)備 aclInit(nullptr); aclrtSetDevice(0); // 創(chuàng)建上下文 aclrtContext context; aclrtCreateContext(context, 0); // 加載模型 aclmdlDesc *modelDesc aclmdlCreateDesc(); aclmdlLoadFromFile(model.om, modelId); aclmdlGetDesc(modelDesc, modelId); // 準(zhǔn)備輸入輸出內(nèi)存 void *inputBuffer nullptr; void *outputBuffer nullptr; aclmdlSizeOfBuffer(modelDesc, ACL_INPUT); aclrtMalloc(inputBuffer, inputSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(outputBuffer, outputSize, ACL_MEM_MALLOC_HUGE_FIRST); // 執(zhí)行推理 aclmdlExecute(modelId, inputBuffer, outputBuffer); // 釋放資源 aclrtFree(inputBuffer); aclrtFree(outputBuffer); aclmdlUnload(modelId); aclrtDestroyContext(context); aclrtResetDevice(0); aclFinalize();這個(gè)骨架里需要注意的細(xì)節(jié)是輸入輸出內(nèi)存的對(duì)齊和生命周期。推理API要求輸入輸出內(nèi)存必須是本設(shè)備可訪問的內(nèi)存由特定接口分配而不是普通malloc來的內(nèi)存。如果在PC端開發(fā)調(diào)試直接用到PC內(nèi)存會(huì)導(dǎo)致推理失敗。這塊在文檔里寫得清楚但很多人第一遍跑不通大多是栽在這里。6.3 多路并發(fā)推理時(shí)的兩個(gè)隱蔽問題端側(cè)場(chǎng)景往往不是單次推理而是要同時(shí)處理多路請(qǐng)求。多路并發(fā)時(shí)有兩個(gè)問題特別隱蔽。第一個(gè)是設(shè)備上下文的線程安全問題。ACL的上下文是非線程安全的每個(gè)線程需要獨(dú)立創(chuàng)建自己的上下文不能多個(gè)線程共享同一個(gè)上下文執(zhí)行推理。曾經(jīng)有一次排查性能問題時(shí)發(fā)現(xiàn)兩個(gè)線程同時(shí)用一個(gè)上下文結(jié)果推理結(jié)果錯(cuò)亂排查了很久才發(fā)現(xiàn)是上下文沖突。解決方法是每個(gè)線程創(chuàng)建獨(dú)立上下文互不干擾。第二個(gè)是Batch推理的輸入拼接問題。如果業(yè)務(wù)場(chǎng)景里每個(gè)請(qǐng)求的數(shù)據(jù)尺寸不一致就不能簡單拼成一個(gè)大Batch送進(jìn)去需要做Padding到統(tǒng)一尺寸或者在模型側(cè)支持動(dòng)態(tài)Batch。Padding到統(tǒng)一尺寸會(huì)浪費(fèi)計(jì)算但實(shí)現(xiàn)簡單動(dòng)態(tài)Batch省算力但配置復(fù)雜。實(shí)際項(xiàng)目里要看端側(cè)設(shè)備的算力余量算力充足就Padding算力緊張就做動(dòng)態(tài)Batch。多路并發(fā)還有一個(gè)容易被忽略的性能點(diǎn)請(qǐng)求調(diào)度順序。如果多個(gè)請(qǐng)求的輸入尺寸差異很大小請(qǐng)求排在大請(qǐng)求后面會(huì)讓小請(qǐng)求等很久影響單次響應(yīng)延遲。一種常用的策略是小請(qǐng)求優(yōu)先大請(qǐng)求稍微靠后用響應(yīng)延遲的均衡換取整體吞吐的穩(wěn)定這屬于業(yè)務(wù)調(diào)度層能做的手腳跟算子優(yōu)化一樣能帶來可觀收益。7. 真實(shí)項(xiàng)目里的調(diào)優(yōu)鏈路從Profile到驗(yàn)證的閉環(huán)7.1 怎么用Profiling數(shù)據(jù)找到真正的熱點(diǎn)而不是靠猜整個(gè)鏈路搭起來之后性能是否達(dá)標(biāo)必須用Profiling工具說話不能靠經(jīng)驗(yàn)猜。Profiling數(shù)據(jù)通常包含幾個(gè)維度算子級(jí)耗時(shí)、數(shù)據(jù)搬運(yùn)耗時(shí)、計(jì)算單元利用率、內(nèi)存帶寬利用率、各核心負(fù)載分布。拿到Profiling結(jié)果后的分析順序我通常按下面幾步來走先看各算子耗時(shí)排名找到Top耗時(shí)算子?？碩op算子的計(jì)算利用率判斷是計(jì)算瓶頸還是搬運(yùn)瓶頸。查看搬運(yùn)耗時(shí)占比如果搬運(yùn)時(shí)間超過計(jì)算時(shí)間優(yōu)先優(yōu)化數(shù)據(jù)復(fù)用和搬運(yùn)連續(xù)化。查看各核心負(fù)載分布如果負(fù)載不均調(diào)整Tiling的多核切分策略。最后才看指令流水線細(xì)節(jié)比如是否因數(shù)據(jù)依賴導(dǎo)致流水線停擺。有一個(gè)真實(shí)案例某模型部署后整體延遲超標(biāo)Profiling顯示一個(gè)卷積算子占了60%耗時(shí)。但我一開始沒看利用率直接優(yōu)化搬運(yùn)邏輯折騰了兩天效果一般。后來重新看數(shù)據(jù)發(fā)現(xiàn)這個(gè)算子計(jì)算利用率只有40%出頭問題出在計(jì)算單元利用率不足——Tiling切得太碎每塊數(shù)據(jù)算的時(shí)間太短搬運(yùn)間隔太長計(jì)算單元大量時(shí)間在等待。于是調(diào)整切塊尺寸讓每塊計(jì)算時(shí)間拉長利用率直接到了70%整體延遲降了一半。這個(gè)案例提醒我Profiling數(shù)據(jù)要看全單看一個(gè)維度容易誤診。算子耗時(shí)高只是癥狀病因可能是計(jì)算、搬運(yùn)、調(diào)度任何一個(gè)環(huán)節(jié)。7.2 性能優(yōu)化常用手段清單按性價(jià)比排序這些年做算子優(yōu)化我把常用手段按性價(jià)比排了個(gè)序新手可以直接照這個(gè)順序來優(yōu)先級(jí)優(yōu)化手段預(yù)期收益實(shí)現(xiàn)成本高完善Tiling切分增大數(shù)據(jù)復(fù)用明顯常能翻倍低高數(shù)據(jù)搬運(yùn)連續(xù)化與對(duì)齊明顯低高啟用雙緩沖或多級(jí)流水線明顯減少等待中中指令選型升級(jí)標(biāo)量→矢量→矩陣部分場(chǎng)景收益大中中融合前后端算子減少落盤明顯但受限中中多核切分負(fù)載均衡視負(fù)載不均程度中低指令級(jí)微調(diào)循環(huán)展開、依賴消除有限高這個(gè)排序不是絕對(duì)的但大方向是對(duì)的先解決數(shù)據(jù)層面的問題再解決計(jì)算層面的問題最后才做指令級(jí)微調(diào)。我見過有人花幾天做循環(huán)展開收益不到3%但其實(shí)他的Tiling切分一塌糊涂數(shù)據(jù)復(fù)用率低得驚人——那種情況下把Tiling改好收益以倍計(jì)。還有一個(gè)基本的驗(yàn)證習(xí)慣每次優(yōu)化只改一個(gè)變量改完就跑Profiling對(duì)比確認(rèn)收益來自這個(gè)改動(dòng)而不是其他巧合。同時(shí)優(yōu)化多個(gè)變量很難定位到底是誰起的效果也會(huì)讓后續(xù)回歸測(cè)試變得不可控。7.3 端側(cè)部署后的驗(yàn)證與回歸性能達(dá)標(biāo)不算完性能調(diào)優(yōu)完成后還有一個(gè)必須做的環(huán)節(jié)功能正確性驗(yàn)證和精度回歸。算子的數(shù)值計(jì)算邏輯和框架原始實(shí)現(xiàn)的輸出可能存在微小差異尤其是使用混合精度或快速算法后誤差會(huì)變大。端側(cè)場(chǎng)景里如果對(duì)精度有硬性要求比如分類任務(wù)的Top-1準(zhǔn)確率需要在部署前后做對(duì)比測(cè)試。我習(xí)慣搭建一個(gè)對(duì)比測(cè)試腳本輸入同一組測(cè)試數(shù)據(jù)分別用原始框架在PC端跑和轉(zhuǎn)換后的模型在端側(cè)跑逐層對(duì)比輸出。對(duì)比時(shí)用兩個(gè)指標(biāo)絕對(duì)誤差和相對(duì)誤差設(shè)置一個(gè)容忍閾值。如果某些層的誤差超過閾值優(yōu)先檢查這些層的算子實(shí)現(xiàn)是否存在數(shù)值穩(wěn)定性問題。還有一個(gè)部署后容易忽略的回歸項(xiàng)連續(xù)運(yùn)行穩(wěn)定性。端側(cè)設(shè)備長時(shí)間推理可能出現(xiàn)內(nèi)存泄漏、性能衰減、溫度過高導(dǎo)致的降頻等問題。建議壓測(cè)至少持續(xù)幾小時(shí)觀察延遲曲線是否平穩(wěn)。我碰到過一次推理延遲初期正常一小時(shí)后逐漸變高最后定位是推理上下文沒釋放導(dǎo)致內(nèi)存碎片積累內(nèi)存分配越來越慢。這種問題屬于運(yùn)行時(shí)資源管理層面的在Profiling單算子階段根本不會(huì)暴露只能在長時(shí)間壓測(cè)里發(fā)現(xiàn)。最后分享一點(diǎn)個(gè)人體會(huì)全鏈路加速這件事沒有一勞永逸的解決方案。算子優(yōu)化、圖編譯、端側(cè)部署、性能調(diào)優(yōu)每個(gè)環(huán)節(jié)單獨(dú)拎出來都有大量細(xì)節(jié)但它們永遠(yuǎn)是互相牽連的。算子寫得好圖優(yōu)化省力圖優(yōu)化做得好部署階段的內(nèi)存焦慮就小部署配置得當(dāng)性能調(diào)優(yōu)才不會(huì)面對(duì)一個(gè)先天不足的模型。我最早做算子優(yōu)化時(shí)總想把每行代碼都調(diào)到極致后來發(fā)現(xiàn)粒度太細(xì)反而失去全局視野?，F(xiàn)在的習(xí)慣是先從整條鏈路最粗的瓶頸入手——數(shù)據(jù)搬運(yùn)、Tiling、融合這些是決定量級(jí)的優(yōu)化確認(rèn)粗粒度優(yōu)化到位后再回頭處理指令級(jí)細(xì)節(jié)。這樣下來整體鏈路穩(wěn)定可控性能也能達(dá)到預(yù)期范圍。希望這篇文章能幫你少走一些我走過的彎路把整個(gè)鏈路理解成一個(gè)系統(tǒng)來看待而不是零散地東改一處、西改一處。 算子優(yōu)化實(shí)戰(zhàn)：從Tiling到端側(cè)部署的全鏈路加速
1. 為什么算子優(yōu)化才是全鏈路加速的命門先聊一個(gè)很多剛接觸異構(gòu)計(jì)算的同行容易忽略的事實(shí)一套加速鏈路里最耗時(shí)間的往往不是模型結(jié)構(gòu)本身而是算子層執(zhí)行效率。模型結(jié)構(gòu)是算法設(shè)計(jì)者定好的骨架算子是把骨架變成真實(shí)算力的肌肉。肌肉乏力骨架再合理也跑不出性能。我最早接觸CANNCompute Architecture for Neural Networks異構(gòu)計(jì)算架構(gòu)時(shí)也踩過這個(gè)認(rèn)知偏差拿到一個(gè)模型第一步就往圖優(yōu)化和部署工具上使勁覺得算子實(shí)現(xiàn)是芯片廠商的事自己寫好算法就行。后來在一次真實(shí)項(xiàng)目里一個(gè)簡單的自定義算子因?yàn)闆]有做Tiling切分端側(cè)推理整條鏈路延遲多了將近一倍才明白算子層是整個(gè)加速流水線的底座。底座不牢后面所有優(yōu)化都是錦上添花成不了雪中送炭。這篇文章圍繞一條完整鏈路來寫從算子的計(jì)算分析、Tiling設(shè)計(jì)、指令選型到圖編譯階段的算子融合策略再到端側(cè)部署時(shí)的模型轉(zhuǎn)換與推理配置最后落到性能調(diào)優(yōu)的閉環(huán)方法。適合已經(jīng)能跑通基礎(chǔ)推理流程、想進(jìn)一步壓榨設(shè)備算力的開發(fā)者也適合剛接觸異構(gòu)編程、想建立全局性能觀的新手。整篇文章不是照搬官方文檔而是把我實(shí)際跑過、調(diào)過、踩過坑的經(jīng)驗(yàn)拆開來講。先明確一個(gè)核心認(rèn)知全鏈路加速不是某一環(huán)的孤軍奮戰(zhàn)而是計(jì)算密度、數(shù)據(jù)搬運(yùn)、內(nèi)存布局、調(diào)度策略四個(gè)維度的協(xié)同。這四件事在算子開發(fā)階段就要想清楚拖到部署階段再補(bǔ)救空間就很有限了。2. 從計(jì)算特征出發(fā)先搞清楚你這算子到底缺什么2.1 算子性能瓶頸的判斷方法拿計(jì)算量和訪存量說話不管下游用的是什么加速芯片算子的性能瓶頸就兩種算力吃滿但數(shù)據(jù)喂不飽叫內(nèi)存瓶頸數(shù)據(jù)搬運(yùn)沒問題但計(jì)算單元在空轉(zhuǎn)叫計(jì)算瓶頸。判斷方法不復(fù)雜先算兩個(gè)數(shù)。以矩陣乘法為例一個(gè) ( M \times K ) 的矩陣乘以 ( K \times N ) 的矩陣計(jì)算量是 ( 2 \times M \times N \times K ) 次浮點(diǎn)操作數(shù)據(jù)搬運(yùn)量是 ( (M \times K K \times N M \times N) \times 4 ) 字節(jié)按FP32算。計(jì)算密度 總計(jì)算量 / 總數(shù)據(jù)量這個(gè)值如果遠(yuǎn)低于設(shè)備的計(jì)算訪存比就是內(nèi)存瓶頸如果高出很多就優(yōu)先優(yōu)化計(jì)算指令的流水線效率。我自己習(xí)慣在動(dòng)手寫算子前先做這個(gè)估算再對(duì)照設(shè)備的理論峰值來判斷優(yōu)化方向。這一步不值錢但能讓優(yōu)化避開南轅北轍的坑。有一個(gè)實(shí)際例子某模型里的歸一化算子輸入是 ( 2048 \times 2048 ) 的矩陣逐元素操作計(jì)算量很小數(shù)據(jù)量很大典型的內(nèi)存瓶頸。這種情況下拼命優(yōu)化計(jì)算指令根本沒用關(guān)鍵是減少重復(fù)搬運(yùn)讓數(shù)據(jù)在片上多復(fù)用。2.2 算子的計(jì)算模式分類決定了它適合走哪條路異構(gòu)芯片雖然指令集各有不同但算子按計(jì)算模式大致能分成幾類分類的意義在于決定優(yōu)化策略的起點(diǎn)逐元素類算子數(shù)據(jù)每個(gè)點(diǎn)獨(dú)立運(yùn)算比如激活函數(shù)、歸一化、逐元素加減。這類算子好寫但極易被內(nèi)存帶寬卡死優(yōu)化重點(diǎn)在數(shù)據(jù)搬運(yùn)和并行切分。歸約類算子需要跨行或跨列聚合比如求和、求均值、Softmax。這類算子的問題在于數(shù)據(jù)依賴需要處理好片上累加和分段歸約。卷積/矩陣乘類算子計(jì)算密度高數(shù)據(jù)復(fù)用性強(qiáng)是最容易榨出峰值算力的類型但Tiling和搬運(yùn)策略也最復(fù)雜。訪存密集類算子比如Transpose、Reshape本身計(jì)算量極小純粹搬運(yùn)優(yōu)化重點(diǎn)是讓搬運(yùn)模式連續(xù)化避免小碎片傳輸。我遇到很多開發(fā)者拿到一個(gè)陌生算子不看計(jì)算特征就先套模板寫寫完了再回頭調(diào)性能效率非常低。正確順序是反過來的先分類再定策略最后才動(dòng)手寫內(nèi)核代碼。確定策略的過程聽起來抽象其實(shí)真正操作時(shí)只需要回答三個(gè)問題數(shù)據(jù)怎么切每一塊放哪一級(jí)緩存計(jì)算用哪種指令完成這三個(gè)問題就是整個(gè)算子開發(fā)階段的核心也是接下來要展開的Tiling設(shè)計(jì)。3. Tiling設(shè)計(jì)把大矩陣切成剛好能塞進(jìn)片上內(nèi)存的形狀3.1 為什么Tiling是算子性能的第一道分水嶺異構(gòu)芯片的存儲(chǔ)層級(jí)和CPU不太一樣通常有全局內(nèi)存相當(dāng)于系統(tǒng)內(nèi)存、片上緩存相當(dāng)于L1/L2、以及緊挨計(jì)算單元的緩沖區(qū)。片上內(nèi)存容量很有限大矩陣不可能一次全放進(jìn)去必須切塊搬運(yùn)算完一塊再搬下一塊。這個(gè)切分策略就是Tiling。Tiling切得好不好直接決定兩層?xùn)|西一是數(shù)據(jù)搬運(yùn)次數(shù)切得太大放不下切得太小搬運(yùn)次數(shù)爆炸二是計(jì)算單元利用率每塊數(shù)據(jù)計(jì)算時(shí)間必須遠(yuǎn)大于搬運(yùn)時(shí)間否則計(jì)算單元大量空轉(zhuǎn)。Tiling里的核心參數(shù)是Tiling Size——在片上內(nèi)存允許的容量內(nèi)盡量讓每一塊算得久一點(diǎn)。比如片上緩沖區(qū)是256KBFP32占4字節(jié)理論上能放65536個(gè)數(shù)。如果做 ( 128 \times 128 ) 的矩陣乘一塊剛好能放下。有兩個(gè)實(shí)際的切分策略固定分塊和自適應(yīng)分塊。固定分塊適合形狀穩(wěn)定的模型代碼簡單自適應(yīng)分塊會(huì)根據(jù)輸入尺寸動(dòng)態(tài)算切法適合輸入形狀變化頻繁的場(chǎng)景。3.2 一個(gè)具體案例512×512矩陣乘的Tiling推導(dǎo)過程光說理論容易飄拿實(shí)際例子算一遍就清楚了。假設(shè)要做 ( C A \times B )A是 ( 512 \times 512 )B也是 ( 512 \times 512 )結(jié)果C同樣是 ( 512 \times 512 )。片上緩沖區(qū)我這里指的是緊挨計(jì)算單元的那個(gè)高速緩沖區(qū)假設(shè)總?cè)萘渴?28KB也就是32768個(gè)FP32數(shù)。如果直接整塊計(jì)算A要占262144個(gè)數(shù)顯然放不下。這時(shí)候按行切分把A按行切成若干塊比如每次取 ( 128 \times 512 ) 的子塊占65536個(gè)數(shù)還是超過32768。繼續(xù)調(diào)整取 ( 64 \times 512 ) 的子塊占32768個(gè)數(shù)剛好卡在容量上限。B不動(dòng)C按行累加。這樣每次搬運(yùn)A的子塊 ( 64 \times 512 32768 ) 個(gè)數(shù)128KBB還是整塊 ( 512 \times 512 262144 ) 個(gè)數(shù)1MB。B在一次算子執(zhí)行里要被反復(fù)讀取8次如果B也太大還要做列切分。從Tiling切分維度來看行切分屬于第一維切分可能還需要第二維切分來降低中間結(jié)果的存儲(chǔ)壓力。這里的關(guān)鍵是把搬運(yùn)量和復(fù)用次數(shù)放到一起看A每次只搬一小條復(fù)用得越充分搬運(yùn)總量越小B必須整體搬屬于不可少的固定開銷。實(shí)踐里我習(xí)慣每塊數(shù)據(jù)在片上至少復(fù)用幾十次才覺得劃算如果復(fù)用次數(shù)太低往往意味著切塊太小浪費(fèi)了片上存儲(chǔ)。3.3 多核切分與負(fù)載均衡別讓一個(gè)核閑著等別人除了片上內(nèi)存維度的切分還有多核維度的切分。設(shè)備上通常有多個(gè)計(jì)算核心每個(gè)核心負(fù)責(zé)一部分?jǐn)?shù)據(jù)。這里最常踩的坑是負(fù)載不均有的核數(shù)據(jù)多、算得慢整個(gè)算子等著它有的核數(shù)據(jù)少、早早算完干瞪眼。避免負(fù)載不均的核心原則是讓每個(gè)核拿到的工作量盡量一致。對(duì)于形狀規(guī)整的矩陣乘按行均分通常就夠了對(duì)于元素稀疏的算子比如某些稀疏計(jì)算數(shù)據(jù)量差異大就需要按非零元素?cái)?shù)量來分配。多核協(xié)作時(shí)的同步設(shè)計(jì)也很關(guān)鍵核與核之間需要通信時(shí)一般走全局內(nèi)存做中轉(zhuǎn)或者用芯片提供的同步機(jī)制。這里需要特別留意不要在大循環(huán)內(nèi)部頻繁同步同步本身就是開銷次數(shù)越少越好。我做過一次實(shí)驗(yàn)同樣的算子把循環(huán)內(nèi)的同步次數(shù)從8次降到2次整體性能提升了大約15%。這種收益不需要改計(jì)算邏輯只需要調(diào)整循環(huán)結(jié)構(gòu)讓每次同步之間有更多獨(dú)立計(jì)算。4. 指令選擇與數(shù)據(jù)搬運(yùn)算得快不如算得巧4.1 計(jì)算指令和搬運(yùn)指令怎么配合才能不打空轉(zhuǎn)Tiling把數(shù)據(jù)切好之后接下來就是真正的指令執(zhí)行。CANN體系下的AI Core編程計(jì)算指令大致分為三種類型矢量指令Vector處理逐元素運(yùn)算面向向量化計(jì)算矩陣指令Matrix/Cube處理矩陣乘加操作計(jì)算密度最高標(biāo)量指令Scalar處理地址計(jì)算、循環(huán)控制等輔助邏輯。真實(shí)項(xiàng)目里寫算子時(shí)最容易忽略的是標(biāo)量指令對(duì)性能的影響。我曾經(jīng)優(yōu)化過一個(gè)算子計(jì)算邏輯本身已經(jīng)用矩陣指令寫得很高效但控制流里用了大量標(biāo)量指令做地址計(jì)算結(jié)果標(biāo)量單元成了瓶頸整體耗時(shí)反而比優(yōu)化前更差。后來把地址計(jì)算提前到搬運(yùn)階段做好緩存循環(huán)內(nèi)只做查表性能才上來。結(jié)論是主循環(huán)里盡量用計(jì)算類指令少放標(biāo)量指令。指令選擇的核心原則是能用矩陣指令解決的不要用矢量指令能用矢量指令解決的不要用標(biāo)量指令一條條算。三種指令之間不是替代關(guān)系而是流水線關(guān)系——標(biāo)量做指揮矢量做批量矩陣做重火力。寫出高效算子的人往往是讓這三者在時(shí)間上重疊數(shù)據(jù)搬運(yùn)的同時(shí)上一塊數(shù)據(jù)正在計(jì)算這一塊數(shù)據(jù)正在準(zhǔn)備搬運(yùn)。4.2 數(shù)據(jù)搬運(yùn)的隱藏成本Copy是性能刺客所有異構(gòu)計(jì)算里數(shù)據(jù)搬運(yùn)的代價(jià)都被嚴(yán)重低估。搬運(yùn)指令本身執(zhí)行不慢慢的是搬運(yùn)前后的等待和地址對(duì)齊問題。我習(xí)慣把數(shù)據(jù)搬運(yùn)分成兩種塊搬運(yùn)整塊連續(xù)數(shù)據(jù)和跳變搬運(yùn)按一定間隔取數(shù)。塊搬運(yùn)效率遠(yuǎn)高于跳變搬運(yùn)這個(gè)差距在某些硬件上能到3倍以上。所以優(yōu)化搬運(yùn)的第一準(zhǔn)則就是讓搬運(yùn)盡量連續(xù)。如果輸入數(shù)據(jù)的原始排布不連續(xù)先做一次重排再搬運(yùn)連續(xù)塊往往比直接多次跳變搬運(yùn)更快。這個(gè)先重排再搬運(yùn)的思路有點(diǎn)反直覺因?yàn)槎嗔艘淮巫x寫但實(shí)測(cè)下來很多場(chǎng)景是賺的。原因是連續(xù)搬運(yùn)能充分利用總線帶寬跳變搬運(yùn)每次都浪費(fèi)大量總線位寬。數(shù)據(jù)搬運(yùn)還有對(duì)齊要求。主流加速芯片普遍要求搬運(yùn)的起始地址和長度按特定字節(jié)對(duì)齊通常是32字節(jié)有些甚至是64字節(jié)。不對(duì)齊的搬運(yùn)要么報(bào)錯(cuò)要么被硬件強(qiáng)行拆分性能直接下滑。所以算子開發(fā)階段通常會(huì)有對(duì)齊這個(gè)參數(shù)它不是一個(gè)可配可不配的開關(guān)而是直接影響搬運(yùn)效率的關(guān)鍵條件。4.3 矢量指令寫算子一個(gè)ReLU算子的完整代碼拆解為了把上面的概念落到代碼層面我用Ascend C寫一個(gè)最簡單的ReLU算子展示矢量指令的基本用法和結(jié)構(gòu)#include kernel_operator.h using namespace AscendC; constexpr int32_t BUFFER_NUM 2; // 雙緩沖 class KernelReLU { public: __aicore__ inline KernelReLU() {} __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, int32_t totalLen) { this-totalLen totalLen; this-blockLength totalLen / GetBlockNum(); // 每核處理長度 xGm.SetGlobalBuffer((__gm__ half*)x GetBlockIdx() * blockLength, blockLength); yGm.SetGlobalBuffer((__gm__ half*)y GetBlockIdx() * blockLength, blockLength); pipe.InitBuffer(inQueue, BUFFER_NUM, blockLength * sizeof(half)); pipe.InitBuffer(outQueue, BUFFER_NUM, blockLength * sizeof(half)); } __aicore__ inline void Process() { CopyIn(); Compute(); CopyOut(); } private: __aicore__ inline void CopyIn() { LocalTensorhalf inLocal inQueue.AllocTensorhalf(); DataCopy(inLocal, xGm, blockLength); inQueue.EnQue(inLocal); } __aicore__ inline void Compute() { LocalTensorhalf inLocal inQueue.DeQuehalf(); LocalTensorhalf outLocal outQueue.AllocTensorhalf(); Maximum(outLocal, inLocal, half(0), blockLength); // 矢量指令取最大值等效ReLU outQueue.EnQue(outLocal); inQueue.FreeTensor(inLocal); } __aicore__ inline void CopyOut() { LocalTensorhalf outLocal outQueue.DeQuehalf(); DataCopy(yGm, outLocal, blockLength); outQueue.FreeTensor(outLocal); } private: TPipe pipe; TGlobalBufferhalf xGm, yGm; TQueQuePosition::VECIN, BUFFER_NUM inQueue; TQueQuePosition::VECOUT, BUFFER_NUM outQueue; int32_t totalLen 0; int32_t blockLength 0; }; extern C __global__ __aicore__ void ReLU(GM_ADDR x, GM_ADDR y, int32_t totalLen) { KernelReLU op; op.Init(x, y, totalLen); op.Process(); }這段代碼里已經(jīng)包含了雙緩沖結(jié)構(gòu)inQueue和outQueue各自有2個(gè)緩沖區(qū)。處理當(dāng)前數(shù)據(jù)塊的同時(shí)下一個(gè)數(shù)據(jù)塊已經(jīng)開始搬運(yùn)搬運(yùn)和計(jì)算重疊起來隱藏了大部分搬運(yùn)延遲。Maximum指令一次處理半精度浮點(diǎn)數(shù)直接走矢量單元語義上等價(jià)于ReLU因?yàn)樗龅氖虑榫褪前沿?fù)值壓到0。這個(gè)例子里有個(gè)細(xì)節(jié)值得留意緩沖區(qū)數(shù)量設(shè)成2是實(shí)現(xiàn)流水線重疊的最基本配置。如果設(shè)成1搬運(yùn)和計(jì)算就必須嚴(yán)格串行性能至少差一半。這是算子開發(fā)里雙緩沖設(shè)計(jì)的關(guān)鍵意義。5. 圖編譯與算子融合讓模型在跑之前先瘦身5.1 編譯期到底干了什么構(gòu)圖、優(yōu)化、內(nèi)存分配三步走算子單獨(dú)寫好了但模型不是一個(gè)個(gè)算子孤立跑的。CANN把整張模型圖交給編譯器統(tǒng)一編排這個(gè)階段通常叫圖編譯。圖編譯的核心工作大致可以分成三步構(gòu)圖與解析、圖優(yōu)化與融合、內(nèi)存分配與調(diào)度。構(gòu)圖階段把模型的結(jié)構(gòu)描述轉(zhuǎn)換成內(nèi)部圖表示這一步會(huì)做算子合法性檢查比如輸入輸出的數(shù)據(jù)類型是否匹配、維度是否符合算子約束。圖優(yōu)化階段是性能的關(guān)鍵包含算子融合、常量折疊、死分支消除等手段。內(nèi)存分配階段決定每個(gè)中間張量放到哪塊物理內(nèi)存、什么時(shí)候釋放做得好的話能大幅降低峰值內(nèi)存占用。我見過不少開發(fā)者繞開圖編譯階段直接把每個(gè)算子單獨(dú)執(zhí)行覺得那樣控制力更強(qiáng)。如果只有一個(gè)算子可能無所謂但在模型級(jí)別這樣做等于放棄了編譯器層面的自動(dòng)優(yōu)化——算子融合、內(nèi)存復(fù)用這些都享受不到了。正確的做法是自定義算子先單獨(dú)驗(yàn)證正確性然后放進(jìn)整個(gè)圖里一起編譯讓編譯器統(tǒng)一編排。5.2 算子融合的具體收益一個(gè)實(shí)測(cè)數(shù)據(jù)算子融合是圖優(yōu)化里收益最直觀的手段。原理不復(fù)雜把多個(gè)可合并成一個(gè)計(jì)算過程的算子融合成一個(gè)算子減少中間張量的落盤和讀回。舉一個(gè)實(shí)際案例。某視覺模型里有一段連續(xù)的卷積、批歸一化、ReLU三個(gè)操作。未融合時(shí)卷積的輸出要寫回全局內(nèi)存批歸一化再從全局內(nèi)存讀出來算完再寫回ReLU再讀一次。三次搬運(yùn)每個(gè)中間結(jié)果都是一整張?zhí)卣鲌D。融合之后卷積計(jì)算的結(jié)果直接留在片上緩沖區(qū)批歸一化和ReLU在這個(gè)結(jié)果上原位計(jì)算最后只寫一次最終結(jié)果。實(shí)測(cè)下來這個(gè)融合把三段串行算子的總耗時(shí)壓縮了約43%主要省的就是中間結(jié)果反復(fù)搬運(yùn)的時(shí)間。這里要特別提一句融合不是所有算子都能做必須滿足數(shù)據(jù)依賴關(guān)系是線性的、中間結(jié)果不需要給其他分支使用的條件。強(qiáng)行融合會(huì)破壞數(shù)據(jù)依賴或增加寄存器壓力得不償失。5.3 自動(dòng)混合精度與內(nèi)存復(fù)用編譯器的隱形加速器除了算子融合圖編譯階段還有兩個(gè)容易被忽視的優(yōu)化手段自動(dòng)混合精度和內(nèi)存復(fù)用。自動(dòng)混合精度是把一部分算子從FP32降到FP16執(zhí)行在幾乎不影響精度的情況下把計(jì)算量減半、搬運(yùn)量減半。但這不是無腦全降某些對(duì)數(shù)值敏感的操作比如Loss計(jì)算、部分歸一化層必須保留FP32。CANN的自動(dòng)混合精度機(jī)制有算子白名單和黑名單策略能自動(dòng)識(shí)別哪些算子適合降精度。內(nèi)存復(fù)用則是把生命周期不重疊的中間張量分配到同一塊物理內(nèi)存。比如兩個(gè)算子的輸出不會(huì)同時(shí)存在它們就可以共享內(nèi)存。這個(gè)優(yōu)化對(duì)端側(cè)部署尤其重要因?yàn)槎藗?cè)內(nèi)存非常有限復(fù)用得好的模型能省出20%到30%的內(nèi)存占用直接關(guān)系到模型能不能在這臺(tái)設(shè)備上跑起來。6. 端側(cè)部署從模型文件到高效推理服務(wù)的完整鏈路6.1 ATC模型轉(zhuǎn)換時(shí)那些容易被忽略的配置項(xiàng)模型在PC端訓(xùn)練好之后不能直接拿到端側(cè)設(shè)備上跑需要經(jīng)過一步模型轉(zhuǎn)換。在CANN體系里這個(gè)轉(zhuǎn)換工具通常叫ATCAscend Tensor Compiler負(fù)責(zé)把訓(xùn)練框架的模型格式轉(zhuǎn)換成端側(cè)推理設(shè)備能執(zhí)行的格式。ATC轉(zhuǎn)換的核心輸入包括原始模型文件、輸入數(shù)據(jù)的形狀和格式定義、輸出節(jié)點(diǎn)名、目標(biāo)設(shè)備型號(hào)。這幾個(gè)信息一個(gè)都不能錯(cuò)。我最常踩的坑是輸入形狀定義不匹配訓(xùn)練時(shí)模型是動(dòng)態(tài)Shape輸入尺寸每次都可能變但轉(zhuǎn)換時(shí)如果不顯式指定Shape范圍轉(zhuǎn)換出來的模型很可能只支持固定尺寸推理時(shí)一旦輸入尺寸變了直接報(bào)錯(cuò)。端側(cè)部署的常見場(chǎng)景里輸入形狀可能是固定大小也可能在一定范圍內(nèi)浮動(dòng)把握好這兩類需求后就能在模板和邊界條件之間做好取舍。如果需要在推理階段動(dòng)態(tài)改變輸入尺寸就得配置動(dòng)態(tài)Shape功能同時(shí)設(shè)置好Shape范圍。范圍設(shè)太大會(huì)浪費(fèi)內(nèi)存和算力范圍設(shè)太小超出范圍的輸入會(huì)被拒絕。合理的做法是根據(jù)真實(shí)業(yè)務(wù)場(chǎng)景統(tǒng)計(jì)輸入尺寸分布覆蓋絕大多數(shù)情況即可沒必要追求理論上的全范圍支持。轉(zhuǎn)換后生成的文件一般包含模型結(jié)構(gòu)和權(quán)重?cái)?shù)據(jù)是端側(cè)推理引擎直接加載的格式。轉(zhuǎn)換過程還會(huì)做一次圖優(yōu)化把前面提到的融合、量化等策略固化到最終模型里所以ATC階段對(duì)推理性能的影響不亞于算子開發(fā)階段。6.2 ACL推理流程的最小可跑骨架模型轉(zhuǎn)換完成后端側(cè)推理一般通過ACLAscendCL推理API來執(zhí)行。這套API的設(shè)計(jì)思路和很多推理運(yùn)行時(shí)類似初始化設(shè)備、加載模型、創(chuàng)建輸入輸出、執(zhí)行推理、釋放資源。最小可跑的推理流程骨架如下// 初始化設(shè)備 aclInit(nullptr); aclrtSetDevice(0); // 創(chuàng)建上下文 aclrtContext context; aclrtCreateContext(context, 0); // 加載模型 aclmdlDesc *modelDesc aclmdlCreateDesc(); aclmdlLoadFromFile(model.om, modelId); aclmdlGetDesc(modelDesc, modelId); // 準(zhǔn)備輸入輸出內(nèi)存 void *inputBuffer nullptr; void *outputBuffer nullptr; aclmdlSizeOfBuffer(modelDesc, ACL_INPUT); aclrtMalloc(inputBuffer, inputSize, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(outputBuffer, outputSize, ACL_MEM_MALLOC_HUGE_FIRST); // 執(zhí)行推理 aclmdlExecute(modelId, inputBuffer, outputBuffer); // 釋放資源 aclrtFree(inputBuffer); aclrtFree(outputBuffer); aclmdlUnload(modelId); aclrtDestroyContext(context); aclrtResetDevice(0); aclFinalize();這個(gè)骨架里需要注意的細(xì)節(jié)是輸入輸出內(nèi)存的對(duì)齊和生命周期。推理API要求輸入輸出內(nèi)存必須是本設(shè)備可訪問的內(nèi)存由特定接口分配而不是普通malloc來的內(nèi)存。如果在PC端開發(fā)調(diào)試直接用到PC內(nèi)存會(huì)導(dǎo)致推理失敗。這塊在文檔里寫得清楚但很多人第一遍跑不通大多是栽在這里。6.3 多路并發(fā)推理時(shí)的兩個(gè)隱蔽問題端側(cè)場(chǎng)景往往不是單次推理而是要同時(shí)處理多路請(qǐng)求。多路并發(fā)時(shí)有兩個(gè)問題特別隱蔽。第一個(gè)是設(shè)備上下文的線程安全問題。ACL的上下文是非線程安全的每個(gè)線程需要獨(dú)立創(chuàng)建自己的上下文不能多個(gè)線程共享同一個(gè)上下文執(zhí)行推理。曾經(jīng)有一次排查性能問題時(shí)發(fā)現(xiàn)兩個(gè)線程同時(shí)用一個(gè)上下文結(jié)果推理結(jié)果錯(cuò)亂排查了很久才發(fā)現(xiàn)是上下文沖突。解決方法是每個(gè)線程創(chuàng)建獨(dú)立上下文互不干擾。第二個(gè)是Batch推理的輸入拼接問題。如果業(yè)務(wù)場(chǎng)景里每個(gè)請(qǐng)求的數(shù)據(jù)尺寸不一致就不能簡單拼成一個(gè)大Batch送進(jìn)去需要做Padding到統(tǒng)一尺寸或者在模型側(cè)支持動(dòng)態(tài)Batch。Padding到統(tǒng)一尺寸會(huì)浪費(fèi)計(jì)算但實(shí)現(xiàn)簡單動(dòng)態(tài)Batch省算力但配置復(fù)雜。實(shí)際項(xiàng)目里要看端側(cè)設(shè)備的算力余量算力充足就Padding算力緊張就做動(dòng)態(tài)Batch。多路并發(fā)還有一個(gè)容易被忽略的性能點(diǎn)請(qǐng)求調(diào)度順序。如果多個(gè)請(qǐng)求的輸入尺寸差異很大小請(qǐng)求排在大請(qǐng)求后面會(huì)讓小請(qǐng)求等很久影響單次響應(yīng)延遲。一種常用的策略是小請(qǐng)求優(yōu)先大請(qǐng)求稍微靠后用響應(yīng)延遲的均衡換取整體吞吐的穩(wěn)定這屬于業(yè)務(wù)調(diào)度層能做的手腳跟算子優(yōu)化一樣能帶來可觀收益。7. 真實(shí)項(xiàng)目里的調(diào)優(yōu)鏈路從Profile到驗(yàn)證的閉環(huán)7.1 怎么用Profiling數(shù)據(jù)找到真正的熱點(diǎn)而不是靠猜整個(gè)鏈路搭起來之后性能是否達(dá)標(biāo)必須用Profiling工具說話不能靠經(jīng)驗(yàn)猜。Profiling數(shù)據(jù)通常包含幾個(gè)維度算子級(jí)耗時(shí)、數(shù)據(jù)搬運(yùn)耗時(shí)、計(jì)算單元利用率、內(nèi)存帶寬利用率、各核心負(fù)載分布。拿到Profiling結(jié)果后的分析順序我通常按下面幾步來走先看各算子耗時(shí)排名找到Top耗時(shí)算子?？碩op算子的計(jì)算利用率判斷是計(jì)算瓶頸還是搬運(yùn)瓶頸。查看搬運(yùn)耗時(shí)占比如果搬運(yùn)時(shí)間超過計(jì)算時(shí)間優(yōu)先優(yōu)化數(shù)據(jù)復(fù)用和搬運(yùn)連續(xù)化。查看各核心負(fù)載分布如果負(fù)載不均調(diào)整Tiling的多核切分策略。最后才看指令流水線細(xì)節(jié)比如是否因數(shù)據(jù)依賴導(dǎo)致流水線停擺。有一個(gè)真實(shí)案例某模型部署后整體延遲超標(biāo)Profiling顯示一個(gè)卷積算子占了60%耗時(shí)。但我一開始沒看利用率直接優(yōu)化搬運(yùn)邏輯折騰了兩天效果一般。后來重新看數(shù)據(jù)發(fā)現(xiàn)這個(gè)算子計(jì)算利用率只有40%出頭問題出在計(jì)算單元利用率不足——Tiling切得太碎每塊數(shù)據(jù)算的時(shí)間太短搬運(yùn)間隔太長計(jì)算單元大量時(shí)間在等待。于是調(diào)整切塊尺寸讓每塊計(jì)算時(shí)間拉長利用率直接到了70%整體延遲降了一半。這個(gè)案例提醒我Profiling數(shù)據(jù)要看全單看一個(gè)維度容易誤診。算子耗時(shí)高只是癥狀病因可能是計(jì)算、搬運(yùn)、調(diào)度任何一個(gè)環(huán)節(jié)。7.2 性能優(yōu)化常用手段清單按性價(jià)比排序這些年做算子優(yōu)化我把常用手段按性價(jià)比排了個(gè)序新手可以直接照這個(gè)順序來優(yōu)先級(jí)優(yōu)化手段預(yù)期收益實(shí)現(xiàn)成本高完善Tiling切分增大數(shù)據(jù)復(fù)用明顯常能翻倍低高數(shù)據(jù)搬運(yùn)連續(xù)化與對(duì)齊明顯低高啟用雙緩沖或多級(jí)流水線明顯減少等待中中指令選型升級(jí)標(biāo)量→矢量→矩陣部分場(chǎng)景收益大中中融合前后端算子減少落盤明顯但受限中中多核切分負(fù)載均衡視負(fù)載不均程度中低指令級(jí)微調(diào)循環(huán)展開、依賴消除有限高這個(gè)排序不是絕對(duì)的但大方向是對(duì)的先解決數(shù)據(jù)層面的問題再解決計(jì)算層面的問題最后才做指令級(jí)微調(diào)。我見過有人花幾天做循環(huán)展開收益不到3%但其實(shí)他的Tiling切分一塌糊涂數(shù)據(jù)復(fù)用率低得驚人——那種情況下把Tiling改好收益以倍計(jì)。還有一個(gè)基本的驗(yàn)證習(xí)慣每次優(yōu)化只改一個(gè)變量改完就跑Profiling對(duì)比確認(rèn)收益來自這個(gè)改動(dòng)而不是其他巧合。同時(shí)優(yōu)化多個(gè)變量很難定位到底是誰起的效果也會(huì)讓后續(xù)回歸測(cè)試變得不可控。7.3 端側(cè)部署后的驗(yàn)證與回歸性能達(dá)標(biāo)不算完性能調(diào)優(yōu)完成后還有一個(gè)必須做的環(huán)節(jié)功能正確性驗(yàn)證和精度回歸。算子的數(shù)值計(jì)算邏輯和框架原始實(shí)現(xiàn)的輸出可能存在微小差異尤其是使用混合精度或快速算法后誤差會(huì)變大。端側(cè)場(chǎng)景里如果對(duì)精度有硬性要求比如分類任務(wù)的Top-1準(zhǔn)確率需要在部署前后做對(duì)比測(cè)試。我習(xí)慣搭建一個(gè)對(duì)比測(cè)試腳本輸入同一組測(cè)試數(shù)據(jù)分別用原始框架在PC端跑和轉(zhuǎn)換后的模型在端側(cè)跑逐層對(duì)比輸出。對(duì)比時(shí)用兩個(gè)指標(biāo)絕對(duì)誤差和相對(duì)誤差設(shè)置一個(gè)容忍閾值。如果某些層的誤差超過閾值優(yōu)先檢查這些層的算子實(shí)現(xiàn)是否存在數(shù)值穩(wěn)定性問題。還有一個(gè)部署后容易忽略的回歸項(xiàng)連續(xù)運(yùn)行穩(wěn)定性。端側(cè)設(shè)備長時(shí)間推理可能出現(xiàn)內(nèi)存泄漏、性能衰減、溫度過高導(dǎo)致的降頻等問題。建議壓測(cè)至少持續(xù)幾小時(shí)觀察延遲曲線是否平穩(wěn)。我碰到過一次推理延遲初期正常一小時(shí)后逐漸變高最后定位是推理上下文沒釋放導(dǎo)致內(nèi)存碎片積累內(nèi)存分配越來越慢。這種問題屬于運(yùn)行時(shí)資源管理層面的在Profiling單算子階段根本不會(huì)暴露只能在長時(shí)間壓測(cè)里發(fā)現(xiàn)。最后分享一點(diǎn)個(gè)人體會(huì)全鏈路加速這件事沒有一勞永逸的解決方案。算子優(yōu)化、圖編譯、端側(cè)部署、性能調(diào)優(yōu)每個(gè)環(huán)節(jié)單獨(dú)拎出來都有大量細(xì)節(jié)但它們永遠(yuǎn)是互相牽連的。算子寫得好圖優(yōu)化省力圖優(yōu)化做得好部署階段的內(nèi)存焦慮就小部署配置得當(dāng)性能調(diào)優(yōu)才不會(huì)面對(duì)一個(gè)先天不足的模型。我最早做算子優(yōu)化時(shí)總想把每行代碼都調(diào)到極致后來發(fā)現(xiàn)粒度太細(xì)反而失去全局視野?，F(xiàn)在的習(xí)慣是先從整條鏈路最粗的瓶頸入手——數(shù)據(jù)搬運(yùn)、Tiling、融合這些是決定量級(jí)的優(yōu)化確認(rèn)粗粒度優(yōu)化到位后再回頭處理指令級(jí)細(xì)節(jié)。這樣下來整體鏈路穩(wěn)定可控性能也能達(dá)到預(yù)期范圍。希望這篇文章能幫你少走一些我走過的彎路把整個(gè)鏈路理解成一個(gè)系統(tǒng)來看待而不是零散地東改一處、西改一處。