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