一顆 AMD Ryzen 5 5500 的單核心，理論上一秒鐘能吐出 134.4 GFLOPS 的浮點運算，但用最直覺的三層 for 迴圈寫矩陣乘法，實測連零頭都拿不到。同一顆核心、同一套 AVX2 指令集，一份開源專案靠 28 組手工調校把數字一路推到 85.30 GFLOPS——理論峰值的 63.5%——沒換晶片，也沒動用第二顆核心。
單核 Zen 3 榨到 85.3 GFLOPS：一場 FP32 GEMM 的極限逼近
這份題為《85.30 GFLOPS Single-Core FP32 Matrix Multiplication on AMD Zen 3》的 GitHub 專案，記錄了在單一 Zen 3 核心上優化 FP32 GEMM（矩陣乘法）的完整過程：從基本的二維分塊起跑，測了 28 組設定，最後鎖定在暫存器裡塞 4 條累加器、把 FMA 指令排成 4 條獨立鏈、再把 B 矩陣即時打包進連續緩衝區的組合，拿下 85.30 GFLOPS。過程沒有一步是「換更快的硬體」，全是把同一顆核心既有的浮點單元、快取階層與暫存器檔案的容量榨乾——這本身就是一堂微架構課，拆開來看每一步到底在解決什麼問題，比只看最終數字有意思得多。整份 README 用一連串編號設定（MX01 到 MX28）逐步記錄每次調整帶來的漲跌，也把哪些方向試過但失敗的原因寫得很清楚——這種連失敗都攤開來看的記錄方式，比單純秀出最終成績更適合拿來理解微架構限制實際卡在哪裡。
切換絕對值／峰值百分比 · 5 個關鍵設定
理論峰值與二維分塊的起跑點
134.4 GFLOPS 這個天花板不是隨便定的，是 Zen 3 這顆核心的浮點單元硬性規格算出來的：兩個獨立的 FMA（fused multiply-add）執行埠，每個埠一次處理 8 個 32 位元浮點數（AVX2 的 256 位元寬向量），FMA 本身一次算兩個運算（乘加各記一次），再乘上 4.2 GHz 的時脈——2×8×2×4.2＝134.4。這是「每個週期兩個埠都在發射滿寬向量指令、完全沒有等待」的理想狀態，任何一段程式碼只要浮點單元有一個週期閒著，實測數字就會往下掉。
README 提到最終的 85.30 GFLOPS「比原始未優化實作快 56.5 倍」；換算回去，那個沒做任何分塊、直接寫三層 for 迴圈的版本大約落在 1.5 GFLOPS 上下——連理論峰值的百分之一都不到。這個 56 倍的落差，正是整篇文章要拆解的謎題：同一顆核心、同一套指令集，效能差距全部來自「資料怎麼被送進浮點單元」，不是浮點單元本身不夠快。多數人隨手寫矩陣乘法時的直覺寫法，其實就落在這個量級——迴圈順序、記憶體存取模式沒特別安排，浮點單元絕大多數時間都在等資料送到，而不是在真正計算。
第一個真正拿出手的設定叫 MX01，README 標成「Bt + 2Dtile + 8acc（baseline）」，35.16 GFLOPS，佔峰值 26.2%。光是把矩陣切成小塊（二維分塊）、搭配 8 條累加器，就已經吃下四分之一的峰值——這說明矩陣乘法最先要解決的瓶頸不是運算量，而是資料重複使用率。原始三層迴圈每算一個 C 元素都要重新把整列 A、整欄 B 掃過一遍，記憶體頻寬遠遠追不上浮點單元的胃口；分塊之後，同一小塊資料在被踢出快取之前能被多次重複使用，浮點單元才有東西可算。這也是為什麼 26.2% 這個數字看起來不高，卻是整條進展曲線裡最關鍵的一步——後面每一項優化都是在這個「資料已經懂得重複使用」的地基上疊加，地基沒打好，暫存器編排、FMA 鏈長調得再精細也沒用。
快取感知的區塊切分
Zen 3 這顆核心的快取分三層：L1 32 KB、L2 512 KB、L3 16 MB，一層比一層大、也一層比一層慢。分塊的核心任務，是把矩陣切成剛好塞得進某一層快取的小方塊，讓 CPU 在把這塊資料踢出去之前盡量多次重複使用，而不是每次都要往下一層甚至往 DRAM 討資料。README 特別點出 K 方向的分塊大小 BK 調到 256 時效果最好，理由是「在不撐爆 L2 的前提下把重複使用率推到最高」——BK 太小，重複使用率不夠；BK 太大，工作集直接溢出快取，數字反而會崩。BK 太小其實也不是沒有代價：分塊次數變多，迴圈本身的額外開銷跟著墊高，重複使用率雖然安全，卻沒有被真正榨乾——256 是這份測試裡「安全」與「榨乾」剛好抵銷的那個點。
這個崩法在失敗設定裡看得到具體數字：MX09 把 I 方向分塊設到 128、K 方向設到 256，算出來的工作集是 128×256×4 位元組＝128 KB，遠遠超過 L1 的 32 KB 容量，效能直接摔到 8.38 GFLOPS——比什麼分塊都沒做的 MX01 還慘。分塊不是「切得越細越好」或「切得越粗越好」，是要精準卡在某一層快取的容量邊界內側，多一點點就前功盡棄——MX09 的 8.38 GFLOPS 甚至輸給完全沒分塊的 MX01，代表分塊這個動作本身如果尺寸算錯，帶來的傷害比不分塊還大，因為程式碼多繞了一層分塊邏輯的開銷，換來的卻是更頻繁的快取失誤。在這個二線累加器（2-line accumulator）的分塊規模下，設定從 MX03 的 39.34 GFLOPS 一路調到 MX13 的 53.58 GFLOPS，靠的是搭配後面會談到的 FMA 鏈長與資料封裝逐步疊加，但天花板就卡在 53.58 附近——要再往上，得換更大的暫存器編排方式。
暫存器裡的四條累加器
暫存器分塊（register blocking）解決的是另一層瓶頸：就算資料已經穩穩待在 L1 裡，每次內積運算還是要從快取把 A、B 的值讀進暫存器才能算，這段讀取本身也要花時間。作法是把 C 小塊的好幾列同時放進 YMM 向量暫存器裡常駐，讓同一筆從快取讀進來的 A 值可以連續拿去跟好幾列 C 的累加器做 FMA，而不是讀一次只用一次。列數越多，每讀一筆 A 或 B 能攤提的運算量越多。這個比例在效能分析裡常被稱為算術強度（arithmetic intensity）：同一筆從記憶體搬進來的資料，能配上多少次浮點運算。前面的分塊解決的是「資料留在哪一層快取」，暫存器分塊解決的是再往下一層——「資料留在快取之後，還要不要重複從快取讀」，兩者疊加才是算術強度真正被推高的原因。
但這條路不是「越多越好」。MX26 把累加器攤開到 8 條，理論上需要同時佔用 64 個 YMM 暫存器來放 A、B、C 的中介值，實際可動用的暫存器數量遠遠不夠，48 個被迫溢位回堆疊——溢位代表這些「應該留在暫存器裡」的值反覆被寫回記憶體、再讀回來，等於把好不容易省下來的記憶體存取又加回去，效能直接腰斬到 12.24 GFLOPS，比很多更早期、更保守的設定還差。另一個相關的失敗案例 MX27 把累加器改成靠 K 方向分塊來延長存活期，同樣的道理：累加器活得越久，同時要保住的暫存器越多，壓力一樣會壓垮效能，跌到 9.13 GFLOPS。4 條累加器之所以成為這份專案測出來的最佳點，是因為它剛好卡在「重複使用率夠高」與「暫存器容量還沒被逼到溢位」的交界——這也是為什麼上圖三層堆疊裡，中間那層會是效能最高的那層，而不是最厚的那層。
FMA 鏈與四週期延遲
拖曳或用方向鍵切換 chain 長度 · 6 個資料點
Zen 3 上 FMA 這個指令本身的延遲是 4 個週期——意思是一個 FMA 從發射到它的結果真正可以被下一個運算拿來用，中間要等 4 個週期。如果程式只維護一條累加器（chain1），每個 FMA 都得等前一個算完才能發射，兩個 FMA 埠雖然都閒置著，卻沒有獨立的工作可以塞進去，實測數字自然貼近下限。README 對這件事的說法很直接：chain4（4 條獨立累加器）足以把這段延遲藏住——四條互不依賴的累加鏈同時在飛，只要有一條的結果還沒回來，另外三條隨時有指令可以填進 FMA 埠的發射槽，兩個埠幾乎每個週期都有事做。這跟前面兩層優化的邏輯是同一套思路的延伸：分塊解決「資料在不在快取」，暫存器分塊解決「資料要不要重複讀」，chain 長度解決的則是「運算之間有沒有依賴」——三層瓶頸分別卡在記憶體階層、暫存器檔案、指令流水線三個不同位置，缺一層，另外兩層調得再好也頂多逼近各自的天花板。
上面的曲線正好畫出這個機制的完整形狀：chain1 到 chain4 一路爬升，40.04 一路衝到 85.30；chain4 之後不進反退，chain5（MX19）掉到 42.07，chain6（MX10）更慘，只剩 14.26。原因不是「鏈變長就變慢」，而是延遲早在 chain4 就已經藏完，再加的累加器不會換來更多平行度，只會多佔一個 YMM 暫存器的位置；超過某個門檻，暫存器數量不夠用，編譯器被迫把部分累加器溢位回堆疊，反覆讀寫記憶體的成本直接吃掉先前靠分塊、封裝省下來的一切。chain4 之所以是整份測試的最佳點，是因為它剛好卡在「延遲被藏完」與「暫存器還沒被榨乾」兩條線的交會處。
把 B 搬進連續記憶體：封裝、對齊與那些失敗的嘗試
拖曳分隔線比較兩種 B 存取模式
左半：B 矩陣依原本 stride 直接存取，記憶體位址跳來跳去
把 B 矩陣即時打包成連續記憶體後，內層迴圈從跳著抓變成順著抓，快取命中率上來——這一步把 79.21 GFLOPS 再推到 85.30。
暫存器與 FMA 鏈解決的是「運算怎麼排」，最後一段瓶頸是「資料怎麼搬」。矩陣乘法內迴圈掃 B 矩陣時，天生就是非連續存取——同一欄的元素在記憶體裡隔著一整列的距離，每跳一步就可能換一條快取行，浪費頻寬也拖累 TLB。README 記錄了三種處理方式的效能差距：只做記憶體對齊、不封裝（MX16，「4 線 + chain4 + 對齊」）拿到 72.58 GFLOPS；把整個 B 矩陣先轉置一遍（MX20，「4 線 + chain4 + B 轉置」）拿到 79.21 GFLOPS，佔峰值 58.9%；即時把當前用到的 B 區塊（BK×BJ）複製進一塊連續緩衝區（MX24，「4 線 + chain4 + B 即時封裝，BK=256」），拿到全場最高的 85.30 GFLOPS，佔峰值 63.5%。三個數字背後是同一個道理：把非連續存取轉成連續存取，但轉換的「範圍」與「時機」決定了代價。整個矩陣先轉置一次會佔用額外的一份記憶體、還會把 L3 塞滿沒必要的資料；即時只封裝當下要用的小區塊，代價小很多，收益卻更高。換句話說，B 轉置解決的是「存取模式」問題，但用了一個比實際需求更大的手段——它把整個矩陣都變成連續的，但下一個 K 分塊真正會用到的，其實只有其中一小條；on-the-fly 封裝把「解決問題的範圍」精準縮小到「當下真正要用的那一塊」，這正是它能再往上多擠出幾個 GFLOPS 的原因。
對齊本身也有貢獻：把記憶體用 32 位元組對齊配置，讓載入指令可以用要求對齊的 vmovaps 而不是允許未對齊、但通常慢一點的變體，README 記錄這一項單獨帶來大約 5% 的提升。這些調整疊加起來，才把 MX24 推過 MX20，成為整份測試裡的最終冠軍。
點擊任一張卡片看失敗原因 · 4 個設定
四個看似合理、實測卻更差的設定
點擊上方任一設定
MX21 · 責任邊界
非暫存性 store（_mm256_stream_ps）本來是為「寫入後不會再讀」的資料設計，寫入時故意繞過快取。但矩陣乘法的 C 是讀取－修改－寫回：每次累加都要先讀再寫，用了非暫存性 store 之後每次寫入都讓那條快取行失效，下一次讀取又得從 DRAM 重新載入——把設計目的和實際存取模式對反了，效能崩到 1.24 GFLOPS。
MX26 · 責任邊界
8 條累加器理論上能攤提更多 A、B 讀取，但需要同時佔用 64 個 YMM 暫存器，實際可用的數量遠遠不夠，48 個被迫溢位。溢位代表反覆寫回、讀回記憶體，把分塊與封裝辛苦省下的頻寬又加了回去。
MX27 · 責任邊界
另一條路是靠 K 方向分塊讓累加器在暫存器裡活得更久，但存活期越長，同時要保住的暫存器越多，跟 8 條累加器踩到同一類瓶頸，只是換了個角度撞牆。
MX09 · 責任邊界
BI=128、BK=256 算出來的工作集是 128 KB，超過 L1 的 32 KB 容量四倍。分塊的前提是塊要塞得進目標快取層，這裡剛好示範反例：分塊本身沒有錯，但尺寸算錯了，結果比不分塊更差。
另一個違反直覺的結果：README 記錄手動加上 _mm_prefetch 之後效能反而下降了約 8%，理由是「Zen 3 的硬體預取器已經很有效率」——對這種規律的循序掃描模式，硬體本身就能提前把資料抓進快取，軟體再插一道明確的 prefetch 指令只是多花一份指令發射成本，卻沒有額外收益。分塊、暫存器編排、FMA 鏈長、資料封裝，每一項都是「主動介入才有幫助」；prefetch 在這裡是唯一一項「介入反而扣分」的技巧，剛好提醒一件事：優化不是把所有教科書招式都用上，而是先搞清楚瓶頸實際卡在哪一層，硬體已經處理好的地方不需要再插手。最後的驗證也值得一提：README 記錄前三名（MX24、MX22、MX23）跟未優化的參考實作比對，最大絕對誤差是 0.0，也就是位元級完全一致；其餘設定則有大約 1e-6 等級的 FP32 捨入誤差，落在預期範圍內。這代表 chain4、封裝這類看似激進的重排並沒有犧牲數值正確性，速度上去了，答案還是同一個答案。這一點對任何要把手寫 kernel 用在正式流程裡的人都重要：分塊、暫存器編排、chain 長度調整全部只改變運算「發生的順序」與「暫存的位置」，不改變每個乘加本身算的內容，所以只要不誤用非暫存性 store 這類真正會影響資料一致性的指令，數值結果理論上就該維持一致，README 的比對結果剛好印證了這一點。
The unlock：把 85.30 GFLOPS 拆開看，會發現它不是單一技巧的功勞，而是分塊、暫存器編排、FMA 延遲隱藏、資料封裝四件事同時卡在各自的甜蜜點——任何一項調過頭（分塊太大、累加器太多、chain 太長、轉置整個矩陣）都會把前面幾項的收益吃掉，這正是為什麼作者要測滿 28 組設定才找到那個交會點。理解這套組合拳，比背下 85.30 這個數字本身更有用：換一顆核心、換一套指令集，峰值會變，但「先讓資料留在快取、再讓資料留在暫存器、再把運算的延遲藏起來、最後把存取模式拉直」這個推理順序不會變。