做底層算子優(yōu)化的人大概都繞不開這樣一個念頭明明調(diào)一個庫就能拿到近乎峰值性能的矩陣乘法為什么還要自己去寫一個“DeepGEMM”我最早接手這個項目時也是這么想的直到在某個推理引擎里連續(xù)碰到幾個通用庫搞不定的場景——batch只有1、K特別大、還要把激活融合進(jìn)來我才真正下定決心自己從零搓一個GEMM內(nèi)核。DeepGEMM不是一個復(fù)雜到看不懂的項目恰恰相反它的核心思想非常樸素圍繞特定硬件和特定形狀把通用矩陣乘法GEMM的性能榨到極限。這篇文章不打算講那些花哨的推導(dǎo)而是把我從零實(shí)現(xiàn)到逐步調(diào)優(yōu)的整套思路拆開包括分塊策略、共享內(nèi)存布局、指令調(diào)度、數(shù)值精度和真實(shí)調(diào)參記錄。適合正在做算子開發(fā)、推理引擎部署、模型性能優(yōu)化的人參考哪怕你之前只寫過naive的矩陣乘法跟著這套思路也能寫出可用的高性能版本。1. 為什么還要再寫一個GEMM放棄現(xiàn)成庫的理由1.1 庫調(diào)用掩蓋的三個真相很多人覺得高性能矩陣乘法直接調(diào)庫不就行了自己寫大概率寫不過。這話一半對一半不對。通用庫確實(shí)在“通用”這個尺度上做了極致的優(yōu)化但它優(yōu)化的是統(tǒng)計意義上最常見的形狀、最普適的內(nèi)存布局、最保守的數(shù)值策略。而你一旦進(jìn)入真實(shí)業(yè)務(wù)場景這些“優(yōu)化前提”可能全都不成立。第一個被掩蓋的真相是通用庫對形狀并不敏感。它內(nèi)部會做啟發(fā)式選擇碰到M1、N4096、K4096這樣的推理場景大概率會走一條為“大M大N大K訓(xùn)練形狀”設(shè)計的路徑結(jié)果就是明明只有一兩個block在干活整個設(shè)備空閑了一大半。第二個真相是庫接口封裝了內(nèi)存布局但它假設(shè)你的數(shù)據(jù)是連續(xù)標(biāo)準(zhǔn)布局。一旦你的A矩陣來自某個量化算子輸出、B矩陣是轉(zhuǎn)置后拼接的權(quán)重你用庫時往往得先做一次內(nèi)存重排這個拷貝開銷可能直接把計算帶來的收益吃掉。第三個真相是庫不給你融合機(jī)會。激活函數(shù)、縮放因子、偏置、Clamp這些在推理里基本必然存在但放在通用庫里就成了多次讀寫的開銷。我做DeepGEMM的動機(jī)說白了就三個字要可控。我要能在內(nèi)層循環(huán)里順手把LeakyReLU做了能在K維分塊時動態(tài)加縮放能針對長K小batch的形狀單獨(dú)走一條快速路徑。這些需求不是通用庫能優(yōu)雅滿足的。1.2 通用最優(yōu)不等于場景最優(yōu)這里有個很容易被忽略的點(diǎn)同樣一個GEMM“通用庫最優(yōu)”和“你的場景最優(yōu)”經(jīng)常不是同一種實(shí)現(xiàn)。直接調(diào)庫跑起來可能已經(jīng)很快了比如達(dá)到了理論峰值的70%你覺得自己沒必要再折騰。但注意通用庫為了平衡所有形狀會在內(nèi)核選擇上做大量分支它在你的特定形狀上可能只有60%或65%的效率而那丟失的百分之幾到十幾在端到端推理里往往就是批處理吞吐量的差距。我在DeepGEMM里做的第一個決策就是放棄“什么形狀都最優(yōu)”的幻想明確寫死幾類目標(biāo)場景訓(xùn)練場景下M通常較大比如512到4096推理場景下M可能只有1到64但N和K往往很大。這兩類場景我會走不同的主循環(huán)和分塊參數(shù)而不是硬套同一個kernel。這個決策一開始看起來很“笨”但實(shí)測下來它帶來的收益比后面任何一個指令級優(yōu)化都大。正是因?yàn)榉艞壛送ㄓ眯晕也拍茚槍1的場景把整個C矩陣切成一條長條讓每個線程塊沿著K方向順序掃避免在M維度上產(chǎn)生無效block和同步開銷。而通用庫不敢這么做因?yàn)樗扰袛郙1是否值得專門優(yōu)化判斷邏輯本身也會引入開銷。1.3 DeepGEMM的定位平衡點(diǎn)在哪所以DeepGEMM并不是要取代現(xiàn)成的高性能庫它更像是一把手術(shù)刀針對特定形狀和融合需求做定制。我在項目里保留了naive版本作為正確性基線也保留了面向大訓(xùn)練形狀的“通用快速路徑”再疊加一個面向小batch長K的“推理專用路徑”。這三層代碼共用同一套分塊框架但關(guān)鍵參數(shù)和循環(huán)結(jié)構(gòu)不同。這樣的結(jié)構(gòu)還有一個隱性好處調(diào)試定位變得容易。當(dāng)性能或精度出問題時我不用在幾千行匯編式內(nèi)核里大海撈針而是先判斷當(dāng)前形狀走的是哪條路徑再縮小到對應(yīng)路徑的邏輯。如果你也想動手寫自己的DeepGEMM建議從一開始就保持這種“快速路徑通用路徑”的分層別急著把所有東西揉進(jìn)一個kernel。2. 從計算強(qiáng)度反推設(shè)計先算賬再寫代碼2.1 別問快不快先問瓶頸在哪開始寫分塊代碼之前我先做了一道算術(shù)題。GEMM一次乘加需要2×M×N×K次浮點(diǎn)操作而需要搬進(jìn)計算單元的數(shù)據(jù)量是M×N N×K M×KC矩陣讀改寫按分塊局部化后可以壓縮這里先按全量估算。兩者一比就得到計算強(qiáng)度計算強(qiáng)度 ≈ 2MNK / (MN NK MK)這時候結(jié)論非常明顯當(dāng)M、N、K都很大時計算強(qiáng)度很高一個矩陣乘法本質(zhì)上是計算密集型的性能的天花板由FMA吞吐決定而當(dāng)M很小比如推理場景的M1或者K特別小的時候訪存開銷開始占主導(dǎo)你優(yōu)化指令流水線可能遠(yuǎn)不如優(yōu)化內(nèi)存搬運(yùn)來得有效。這個賬必須先算清楚因?yàn)樗苯記Q定了DeepGEMM的主循環(huán)怎么寫。對于大矩陣核心任務(wù)是把數(shù)據(jù)喂給計算單元讓FFMA永不空轉(zhuǎn)對于小矩陣核心任務(wù)變成減少讀入的數(shù)據(jù)次數(shù)和讓C矩陣盡量留在寄存器或共享內(nèi)存里。兩種形態(tài)的代碼長得完全不一樣。2.2 三層分塊寄存器、共享內(nèi)存、L2的接力賽一旦確定瓶頸是計算接下來的問題就是怎么把M×N×K的大循環(huán)拆成能塞進(jìn)硬件結(jié)構(gòu)的小塊。我的做法是三層分塊每一層對應(yīng)一級存儲。第一層是Block Tile。整個C矩陣按128×128的塊劃分每個線程塊負(fù)責(zé)一個塊這個尺寸是給L2緩存和線程塊并行度用的。第二層是Warp Tile。一個128×128的Block Tile交給4個warp處理每個warp拿32×32的子塊。第三層是Thread Tile。每個warp內(nèi)部有若干線程把32×32繼續(xù)切成每個線程負(fù)責(zé)的微塊我常用的配置是每個線程用8×8大小的累加器矩陣。為什么非要三層因?yàn)槊恳患壌鎯Φ娜萘?、帶寬、延遲都不一樣。寄存器最快但容量最小共享內(nèi)存次之L2再次。三層分塊本質(zhì)上是讓數(shù)據(jù)在這三級存儲里逐級接力——A和B的大塊從全局內(nèi)存搬到共享內(nèi)存再由共享內(nèi)存搬進(jìn)寄存器而C的累加結(jié)果盡量在寄存器里累積完最后再一次性寫回。這也是很多高性能GEMM與naive版本最本質(zhì)的區(qū)別naive版本每個線程直接去全局內(nèi)存讀A和B每算一次乘加就要發(fā)起一次訪存而高性能版本是把數(shù)據(jù)先“囤積”在靠近計算單元的存儲里一次搬運(yùn)、多次復(fù)用。2.3 塊形狀和尺寸的試錯起點(diǎn)很多第一次寫GEMM的人會到處抄別人的Tile參數(shù)抄完發(fā)現(xiàn)性能稀爛然后懷疑是自己機(jī)器不行。其實(shí)尺寸這個東西非常依賴具體硬件的手寫特性別人調(diào)好的參數(shù)未必適合你的場景和硬件代次。我的建議是找一組合理的起點(diǎn)然后做參數(shù)掃描不要憑感覺一錘定音。我最初用的一組參數(shù)是這樣層級尺寸說明Block Tile (BM×BN)128×128照顧L2容量和塊間并行度K維單次迭代 (BK)8或16決定共享內(nèi)存單次搬運(yùn)量Warp Tile (WM×WN)32×32每個warp負(fù)責(zé)的C面積Thread Tile (TM×TN)8×8每線程持有的累加器個數(shù)Warp個數(shù)4一個Block內(nèi)4個warp這個起點(diǎn)不是拍腦袋而是基于硬件資源算出來的每個線程8×8需要64個累加寄存器加上A、B的片段寄存器、地址計算、臨時變量寄存器用量在150到200之間不會爆掉寄存器文件又有足夠的ILP指令級并行隱藏延遲。后續(xù)掃描我一般從“BM/BN是否變大”“BK是否太小導(dǎo)致搬運(yùn)頻繁”兩個方向展開這個我放到后面調(diào)優(yōu)章節(jié)細(xì)說。3. 內(nèi)核循環(huán)的主體結(jié)構(gòu)K維主序與微內(nèi)核展開3.1 主循環(huán)里每個Stage在做什么DeepGEMM的主循環(huán)是沿著K維推進(jìn)的不是沿著M或N維。原因很簡單C矩陣的每一個點(diǎn)都是A的某一行和B的某一列的K維點(diǎn)積沿K方向一次推進(jìn)就能同時更新當(dāng)前Block Tile內(nèi)所有C元素這個特性讓K維主序天然適合做數(shù)據(jù)流水。主循環(huán)里每個Stage做四件事第一步從全局內(nèi)存加載A和B各自的一段到共享內(nèi)存記為階段P0第二步等待共享內(nèi)存就緒這叫同步點(diǎn)第三步每個線程從共享內(nèi)存讀入自己需要的A片段和B片段到寄存器階段P1最后執(zhí)行一批FMA運(yùn)算把這組數(shù)據(jù)乘累加到自己的Thread Tile上階段P2。整個結(jié)構(gòu)用偽代碼表達(dá)大致是P0: load A_tile[by][kk] - smem_A load B_tile[kk][bx] - smem_B __syncthreads() P1: load sA[ty][tx] - reg_A load sB[ty][tx] - reg_B // 按tile坐標(biāo)取數(shù)存在swizzle時需做索引映射 P2: for i in 0..TM-1: for j in 0..TN-1: acc[i][j] reg_A[i] * reg_B[j] 前進(jìn)kk回到P0這個循環(huán)寫好之后我第一時間就不是看性能而是先拿小矩陣做數(shù)值對比。因?yàn)镚EMM一旦分塊最容易出錯的就是索引坐標(biāo)映射一個分塊內(nèi)部的對齊錯位可能只在特定形狀下觸發(fā)非常陰險。3.2 雙緩沖與組播把等待變成計算P0到P2的串行結(jié)構(gòu)有個明顯問題加載共享內(nèi)存是慢的而計算是快的如果每次都要等共享內(nèi)存加載完才開始算計算單元就會頻繁空轉(zhuǎn)。解決辦法是雙緩沖。我在共享內(nèi)存里給A和B各準(zhǔn)備兩塊buffer當(dāng)前正在計算的那塊記為當(dāng)前buffer下一輪要用的數(shù)據(jù)提前加載到另一塊buffer。這樣P2的計算可以和下一次P0的加載重疊GPU在做乘加的同時DMA單元已經(jīng)在下一次數(shù)據(jù)了。同步點(diǎn)也從一個變成兩個一個確保上一輪計算結(jié)束前不能覆蓋當(dāng)前bufferconsumer同步一個確保下次數(shù)據(jù)寫完后才能開始下一輪計算producer同步。組播multicast或者說廣播讀取是另一個關(guān)鍵。同一個warp里的線程在K維推進(jìn)時A的某一行會被多個線程重復(fù)讀取B的某一列也是。與其讓每個線程都去共享內(nèi)存讀一遍不如讓硬件感知到這是同一條數(shù)據(jù)一次讀出來廣播給多個線程。這個優(yōu)化不在代碼層面顯式出現(xiàn)但你的數(shù)據(jù)布局和索引模式會影響硬件能不能自動識別這種廣播模式。這也是為什么我堅持用規(guī)整的、線性偏移的索引方式而不是隨手做各種奇怪的重排。3.3 微內(nèi)核中的指令級流水主循環(huán)里最內(nèi)層的P2部分我把它稱為微內(nèi)核micro-kernel。這一小段代碼寫得好不好直接決定了最終性能能到峰值的百分之多少。以每線程8×8累加器為例微內(nèi)核要做64個乘累加對應(yīng)64條FFMA指令。這里有個經(jīng)驗(yàn)FFMA指令雖然是一個周期就能發(fā)出一條但如果編譯器生成的指令流是“反復(fù)往同一個寄存器上累加”會產(chǎn)生很長的依賴鏈——下一周期要用的值必須等上一周期的結(jié)果算完流水線就堵住了。所以我寫微內(nèi)核時會讓連續(xù)的幾次FFMA依次作用在8個不同的累加器上把一個64步的長依賴鏈打散成8條獨(dú)立的短鏈。這個思路有點(diǎn)像流水線上一個工人只擰一顆螺絲會卡住整條產(chǎn)線不如讓幾個工人各管一條螺釘線。實(shí)際我通常會讓寄存器里的A片段和B片段交錯排列先算acc[0][0]、acc[0][1]、acc[0][2]再算acc[1][0]……繞一圈回來再輪到acc[0][0]的下一次累加。循環(huán)展開因子選4或8太小依賴鏈壓不住太大會把I-Cache擠爆具體數(shù)值要看指令發(fā)射寬度。4. 訪存優(yōu)化的硬仗共享內(nèi)存沖突與Swizzle4.1 一個反直覺的停頓我的DeepGEMM在第一版完成時性能只比naive快了一倍離預(yù)期差得遠(yuǎn)。我拿著profiling結(jié)果一行行找發(fā)現(xiàn)計算單元利用率不低但是共享內(nèi)存的吞吐指標(biāo)異常高像是有東西在共享內(nèi)存上來回打轉(zhuǎn)。后來才意識到是bank conflict在作祟。共享內(nèi)存按bank組織同一bank同一周期只能服務(wù)一次訪問如果同一warp的多個線程同時訪問同一個bank硬件就得把這些訪問串行化。第一個版本我圖省事A矩陣按行主序直接鋪在共享內(nèi)存里結(jié)果一個warp里的線程去讀相鄰列的時候恰好踩到了相同的bank或bank組訪存效率幾乎腰斬。這個問題的反直覺之處在于從高級語言看每次線程訪問的都是“自己的地址”編譯器也沒有任何報錯但性能就是上不去。如果不熟悉bank機(jī)制很可能反過來懷疑是自己循環(huán)展開寫壞了白白浪費(fèi)好幾天排查時間。4.2 XOR Swizzle解決沖突的直覺解決bank conflict最通用的一招是swizzle。我的做法是對共享內(nèi)存里的數(shù)據(jù)做一次XOR映射讓原本會連續(xù)訪問的地址被打散到不同的bank上。直覺上可以這樣理解共享內(nèi)存有32個bank同一warp有32個線程如果讓地址的排列方式滿足“每個線程訪問的bank號互不相同”沖突就消失了。對于A矩陣的共享內(nèi)存布局我會在索引里加一個XOR操作把行號和列號的一部分做一個異或再決定實(shí)際地址偏移。這個操作是在把全局內(nèi)存數(shù)據(jù)寫入共享內(nèi)存時就做好的計算階段讀取時直接用映射后的地址幾乎不增加額外指令開銷。當(dāng)然swizzle不是萬能藥。如果你的訪問模式本身就比較隨機(jī)XOR反而會引入新的碰頭概率。我用它主要是因?yàn)镚EMM里A片段的訪問模式非常固定——同一個warp讀同一行或同一塊列規(guī)律性強(qiáng)swizzle的效果是立竿見影的。關(guān)鍵是要保證映射是雙射也就是不丟失數(shù)據(jù)這個我會先用小規(guī)模地址枚舉驗(yàn)證一遍。4.3 向量化與對齊LDG.STS的節(jié)奏共享內(nèi)存那邊理順以后下一個瓶頸就出現(xiàn)在全局內(nèi)存到共享內(nèi)存的數(shù)據(jù)搬運(yùn)上。GEMM的搬運(yùn)特點(diǎn)是塊狀、連續(xù)、量大非常適合向量化加載。我在DeepGEMM里盡量讓每一次從全局內(nèi)存讀取都是16字節(jié)粒度也就是float4級別的讀取這樣一次訪存能取回4個浮點(diǎn)數(shù)訪存指令數(shù)少了四分之三。這里有個對齊細(xì)節(jié)向量化讀取要求起始地址至少是16字節(jié)對齊。如果一段數(shù)據(jù)恰好從偏移量4字節(jié)開始我就沒法直接用float4要么手動拆分要么在數(shù)據(jù)布局上填padding。我在整個項目里對A和B的leading dimension做了對齊約束寧可在尾部多填一些無效數(shù)據(jù)也不要讓主路徑出現(xiàn)未對齊訪問。向量化和swizzle要配合著做。先從全局內(nèi)存用float4讀入寄存器再按swizzle后的地址把數(shù)據(jù)寫入共享內(nèi)存這個寫入過程也盡量保持向量化。也就是說寄存器里的4個數(shù)要能連續(xù)寫到共享內(nèi)存的連續(xù)4個位置否則就得拆成4次標(biāo)量寫前面省下的訪存指令又吐回去一部分。5. 數(shù)值精度與邊界情況快但必須算得對5.1 低精度輸入、高精度累加的分寸DeepGEMM面向深度學(xué)習(xí)場景輸入經(jīng)常是FP16或BF16但累加器我會堅持用FP32。原因很簡單K維經(jīng)常是4096甚至更大如果一路用FP16累加舍入誤差會在長點(diǎn)積里不斷累積最終結(jié)果可能在十進(jìn)制第四位就開始飄了。分塊累加本身對精度是有幫助的。因?yàn)槊恳恍K只做8或16次乘加局部累計誤差被限制在小范圍內(nèi)最后再用FP32把各塊的累加結(jié)果合起來。這個操作等價于做了一個樹形歸并的簡化版比從頭到尾一長條點(diǎn)積要穩(wěn)得多。不過要注意如果塊尺寸取得過大比如BK64時一次性累加64項局部誤差還是會偏大我一般把BK控制在8到16之間精度和訪存開銷都能兼顧。5.2 Inf/NaN與異常輸入的處理低精度計算里最容易被忽視的是異常值。FP16和BF16的指數(shù)范圍有限哪怕輸入數(shù)據(jù)本身是合理的FP32數(shù)值轉(zhuǎn)成BF16之后也可能溢出成Inf。一旦Inf進(jìn)入乘累加鏈FFMA會把Inf一路傳播下去最后整個C矩陣的對應(yīng)行列全部變成Inf或NaN看上去像是算法崩了。我在DeepGEMM里加了一個可選的縮放機(jī)制在K維主循環(huán)最開始統(tǒng)計當(dāng)前分塊內(nèi)A和B的最大絕對值如果超過低精度格式能安全表示的范圍就整體乘一個縮放因子再進(jìn)低精度路徑。這個操作讓訓(xùn)練和推理時偶爾出現(xiàn)的異常值不會瞬間污染整塊結(jié)果。代價是多了一次數(shù)據(jù)掃描所以我只在網(wǎng)絡(luò)量化或梯度出現(xiàn)異常時才開啟這個路徑。還有一個很容易踩的坑如果輸入本身含NaNFFMA的傳播行為和標(biāo)準(zhǔn)IEEE語義不完全一樣。有些情況下NaN會被當(dāng)作普通值參與運(yùn)算結(jié)果取決于指令的浮點(diǎn)模式。我的做法是在內(nèi)核層面加一個開關(guān)——如果檢測到NaN輸入走一條保守的標(biāo)量路徑雖然慢但結(jié)果和全FP32版本完全對齊。5.3 非對齊Shape的兜底實(shí)際業(yè)務(wù)里M和N幾乎不可能每次都恰好是128的倍數(shù)。我對非對齊shape的處理策略是主體部分用高性能Tile內(nèi)核尾部剩余的若干行或列用一個專門的小Tile內(nèi)核兜底。小Tile內(nèi)核不追求極端性能只追求不炸寄存器和正確性。選擇閾值時有個經(jīng)驗(yàn)如果剩余部分占整個矩陣面積的比例小于3%尾核隨便寫寫就行主路徑是絕對性能大頭如果比例超過15%就要反思是不是主Tile尺寸選得不合適或者干脆把非對齊維度也納入主路徑用mask處理。mask處理的好處是避免兩套kernel切換帶來的額外同步壞處是predicated FMA會降低主路徑效率。我一般以10%作為切換線。這里還想強(qiáng)調(diào)一個正確性驗(yàn)證的方法非對齊shape最容易暴露索引bug。我的常規(guī)操作是拿一個M130、N66、K17這種邊角料尺寸同時跑naive版本和DeepGEMM版本做數(shù)值對比誤差閾值設(shè)到1e-2以內(nèi)。如果這個形狀能過再上大數(shù)據(jù)集。6. 調(diào)優(yōu)實(shí)測從Profiling到參數(shù)收斂6.1 先用Performance Metric鎖定瓶頸DeepGEMM第一個可用版本跑通后我做的第一件事不是盲調(diào)參數(shù)而是用硬件性能計數(shù)器把三類指標(biāo)拉出來計算單元利用率、共享內(nèi)存吞吐、全局內(nèi)存吞吐。這三者的比例能直接告訴我瓶頸在哪。如果計算單元利用率很高但全局內(nèi)存吞吐也高說明數(shù)據(jù)搬運(yùn)還能和計算重疊得更好如果計算單元利用率不到50%但共享內(nèi)存吞吐已接近上限那就是bank conflict或swizzle沒到位如果指令數(shù)看起來很多但不是FFMA占大頭那可能是地址計算和數(shù)據(jù)搬運(yùn)指令過多該考慮常數(shù)下標(biāo)或減少同步點(diǎn)。拿這些數(shù)據(jù)說話的好處是不會因?yàn)椤澳炒胃膭痈杏X變快了”就自我感覺良好。我經(jīng)常遇到的情況是改了循環(huán)展開因子直覺上指令數(shù)少了實(shí)際因?yàn)榧拇嫫饕绯鰧?dǎo)致perf曲線反而下降這種反直覺問題只能靠計數(shù)器和profiling才能發(fā)現(xiàn)。6.2 一次Tile參數(shù)掃描的完整記錄有一次我把目標(biāo)形狀定為M512、N4096、K4096的訓(xùn)練場景跑了三組參數(shù)結(jié)果很有意思參數(shù)組BM×BNBKWM×WNTM×TN計算單元利用率實(shí)測TFLOPS相對基線A128×128832×328×885%18.2基線B128×1281632×328×888%19.04.4%C256×128864×328×1672%16.0-12%A和B的差異主要在BK上。BK從8翻到16共享內(nèi)存搬運(yùn)次數(shù)減半訪存指令變少利用率提升3個點(diǎn)這個收益是實(shí)實(shí)在在的。C組看似更激進(jìn)的Block Tile實(shí)際因?yàn)槊總€線程的Thread Tile變成8×16后寄存器壓力過大導(dǎo)致occupancy下降計算單元喂不飽數(shù)據(jù)反而倒退了12%。這次掃描讓我明白一個原則參數(shù)之間像木桶的木板單獨(dú)調(diào)高一塊不一定會變好反而可能擠爆另一塊。準(zhǔn)確的做法是一組一組調(diào)每次只動一個變量記錄一版結(jié)果再回滾到最佳點(diǎn)做下一個變量。6.3 最終效果與適用邊界經(jīng)過幾輪掃描DeepGEMM在目標(biāo)形狀上穩(wěn)定到了比最初版本快40%左右換句話說從naive到最終版整體提升了一個數(shù)量級。但要誠實(shí)講這套參數(shù)只在目標(biāo)形狀附近有效。M1的推理場景我會完全換一套參數(shù)Block Tile縮到64×64K維推進(jìn)加大到64主循環(huán)里還額外開了mask路徑因?yàn)檫@個場景下訪存已經(jīng)是主要矛盾計算強(qiáng)度太低用計算密集型的參數(shù)跑就是自找麻煩。這里也順帶提一下DeepGEMM適合什么、不適合什么它適合形狀相對固定、需要融合算子、對數(shù)值行為有特殊要求的場景它不適合形狀極其動態(tài)、每次調(diào)用都要切換kernel的開銷、或者你只是需要一個偶爾跑一次的通用矩陣乘法。項目里我同時保留了naive版本和快速路徑就是尊重這個邊界。我現(xiàn)在養(yǎng)成的習(xí)慣是拿到任何一個GEMM需求先寫一版naive作正確性基線再用三層分塊的框架套一層快速路徑然后永遠(yuǎn)只在profiling數(shù)據(jù)指引下動下一步。每調(diào)一版參數(shù)我就把結(jié)果記錄到一張表里包括相對性能和當(dāng)時的硬件指標(biāo)這樣即使兩周后回來看也能立刻知道哪個方向已經(jīng)試過、哪個方向還有潛力。最后再分享一個小技巧調(diào)優(yōu)時留著所有歷史版本的kernel切換開關(guān)別刪舊代碼否則你永遠(yuǎn)無法確認(rèn)某個性能回退到底是新改動引入的還是硬件狀態(tài)波動導(dǎo)致的。DeepGEMM這個項目做到后面最大的收獲反而不是那百分之幾十的性能而是這套“先算賬、再分層、靠數(shù)據(jù)收尾”的優(yōu)化方法論。