兩天,我想自己寫一個推論引擎去挑戰HaloGen,結果差得遠,但也把 AMD 的底給摸了一遍
環境
AMD Strix Halo:Ryzen AI Max+ 395 + Radeon 8060S(gfx1151 / RDNA3.5)
124G 統一記憶體(CPU/GPU 共用);BIOS carve 只留 512M VRAM,其餘走 GTT
NVMe 資料碟(/mnt/nvme-data:模型、ROCmFPX、llama.cpp fork)
系統與驅動 Fedora + ROCm 7.14(llvm-objdump/hipcc 解剖刀)、amd_iommu=off
LLM推論引擎:halogen-flash 0.6.3版、llama.cpp
起點:一開始只在乎Decode速度體驗
痛點:Prefill的落差太大
事情是這樣開始的——去年底DDR5價格起飛,我趁小主機還沒受影響前在淘寶搞了一台AI Max+ 395,當時只是拿來做本地聊天也挺順暢的,今年開始AI Agent大火,但影響 AI Agent 最大的體驗瓶頸,往往是讀指令、讀輸出、爬文、爬代碼等,是屬於 prefill 速度,反而輸出的比例遠少於輸入的 token 數量。雖然 AMD 能利用內存帶寬在 decode 速度追上了老黃,但 prefill 速度卻遠低於同等級 N 卡,尤其長文本衰退更趨明顯。即使用了增量快取如llama.cpp的prompt-cache 或是 sglang的radix tree 等技術可以緩解這樣的瓶頸,但治標不治本,對於增量的部分就是比較痛苦,尤其在累積較長的上下文後要模型去網上爬文簡直要命。
近期我開始注意到社群裡有個叫 halogen-flash-server 的閉源專案,作者 peonist-ai 利用AI寫了一套針對這顆晶片的自製推論引擎,效能遠超官方 ROCm 路徑。這個發現讓我動了一個念頭:既然有人能透過AI Agent做到,那我能不能也試試看,自己動手改一套推論引擎?
於是找了幾個 AI 工具幫忙(Claude 5 Fable產生計畫書,是 Muse Spark 1.3 負責動手實作,Gemini flash 3.8後期提供了完整的分析計畫跟數據翻譯),開始了這兩天的折騰。
更大的背景:這不是單一台機器的問題
其實會一路查下去,不只是因為 Strix Halo 這台機器,而是因為我另外一張獨立顯卡 RX 9070,拿去跟同世代的 RTX 5070 比,長文本推論(prefill)速度也是輸得很難看——這不是個案,是 AMD RDNA3/RDNA4 這整個世代,在 LLM 推論這個特定負載上,普遍輸給同世代 NVIDIA 老遠的現象。
先把客觀規格攤開來看,不是空泛地說「輸」,是真的輸在哪一塊,我們扣除掉NVFP4的特殊加速格式:
拿 RX 9070 XT 跟 RTX 5070 比,兩者的基礎算力(FP32,一般著色器運算)其實 AMD 還略高一點(約 49 TFLOPS vs 31 TFLOPS,AMD 領先)。但只要看向「張量矩陣運算」(AI/Tensor Core)這一塊,AMD 的 FP16 張量運算約 144.5 TFLOPS,NVIDIA 約 123.5 TFLOPS。
LLM 推論裡的 prefill(處理輸入文字的階段),幾乎全部是矩陣乘法跟注意力運算,吃的正是這塊張量核心,不是一般的著色器算力。AMD 在「一般運算」上不輸甚至略贏,但在「LLM 最依賴的那一種運算」上,或許沒那麼簡單。
除了張量核心規格本身,還有兩個結構性的原因,加起來讓實際使用時的差距比規格表看起來更誇張,所以主要優化的方向也朝這走:
- 搬資料的方式比較笨重:NVIDIA 從幾年前的架構開始,就有一套硬體層級的「背景搬運」機制(cp.async),可以讓資料在運算的同時,自動從主記憶體流向晶片上的高速暫存區,不用佔用運算單元的時間。AMD RDNA 這幾代消費級卡沒有對應的硬體機制,資料要搬運,得老老實實地讓運算單元自己動手去搬,這在處理長文本的大矩陣運算時,會產生大量「等資料到位」的空檔,運算單元沒辦法一直保持忙碌。
- 軟體生態系起步較晚、火力較分散:NVIDIA 那邊的底層函式庫,已經被業界磨了好幾年,各種量化格式、各種模型架構都有現成調校好的版本可以直接用。AMD 這邊的對應工具鏈,相對年輕,而且官方資源優先投入在資料中心等級的晶片上,消費級顯卡跟這種 APU 產品,能分到的官方調校資源少很多——這也是為什麼這兩天我查出來的東西,大部分打雜指令的問題,某種程度上反映的是「官方還沒空好好優化這個特定情境」,不是「這件事做不到」。
- 那RDNA4比RDNA3進步在哪呢?:主要還是遊戲表現,像是光柵引擎與光追單元的效能提升,另外RDNA4起多了支援FP8量化格式,讓RDNA4世代能透過fallback運作MXFP4的量化模型。
搞清楚這個大背景之後,我才決定拿 Strix Halo 這台機器,認真試試看能不能靠自己動手,填補一部分這個差距。
先搞清楚差距有多大
以同一台機器AMD Strix Halo平台下跑分測試對照,我測試的對象是 Qwen3.8-27B(dense 模型架構,混合注意力層)。在 32K context 下,我自己跑的 llama.cpp stock 版本 prefill 速度是 263.5 tok/s,halogen 那邊實測是 570.4 tok/s——差了 2.17 倍。同場加映:
| 檔位 |
llama.cpp stock |
halogen |
倍率 |
| pp8K |
317.8 tok/s |
634.9 tok/s |
2.00x |
| pp32K |
263.5 tok/s |
570.4 tok/s |
2.17x |
| pp128K |
137.6 tok/s |
382.2 tok/s |
2.78x |
注意最後一行:差距隨深度拉大——這條曲線後面會成為關鍵線索。
這裡要澄清一個容易搞混的地方:halogen 官方展示的另一個更誇張的數字(1424 tok/s,約 4 倍差距)是跑在另一個模型 Qwen3.8-Flash-Next 上,那個模型用了 512 專家的 MoE 架構,每個 token 只需要算一小部分參數,運算量本來就少很多,不是同一個擂台,不能直接拿來比。我這次挑戰的,是相對公平的 27B dense 對照。
第一步:先看看官方路徑到底把時間花在哪
沒有亂猜,直接把官方 llama.cpp 跑 Q5_K 量化的核心運算(MMQ)利用SASS組合語言解包,一條一條指令去數。整顆 kernel 1724 條指令裡,真正在做矩陣運算的只有 32 條(約 2%,聚焦熱區則約 7%)。結果發現:
- 剩下六成多(62%)的指令,都在處理量化格式的「拆包」——把壓縮過的權重數字,即時解開、搬到對的位置、換算成能運算的格式
- 暫存器用量 81,是健康線 32 的 2.5 倍,導致 GPU 沒辦法同時排進足夠多的工作,進一步拖慢速度
簡單講:真正幹活的部分被一大堆「打雜」的指令包圍,張量核心根本沒吃飽。單算子實測:ffn_up 形狀只跑出 21.0 TFLOPS,紙面峰值 59 的三成六。
野心勃勃的完整規劃:三層自研引擎架構
看到這個發現,我一開始很樂觀,覺得既然問題找到了,那乾脆不要只是修修補補,直接規劃一套完整的自研引擎架構,三層疊起來:
┌────────────────────────────────────────────────────────┐
│ 排程層:Radix Tree 前綴快取(Prompt Caching) │ ── 命中時直接跳過重算,秒回第一個字
└──────────────────────────┬─────────────────────────────┘
│
┌──────────────────────────▼─────────────────────────────┐
│ 記憶體階層管理:統一記憶體動態換出 / 回收 │ ── 榨乾這台機器的大容量記憶體,長文本不爆
└──────────────────────────┬─────────────────────────────┘
│
┌──────────────────────────▼─────────────────────────────┐
│ 執行底座:量化格式直灌張量核心的專用運算核心 │ ── 幹掉打雜指令,直接餵飽張量運算單元
└────────────────────────────────────────────────────────┘
- 排程層(前綴快取):多輪對話或代碼助理這種場景,常常會有重複的開頭內容(系統提示詞、共用的上下文)。做一個前綴共享機制,命中過的內容直接跳過重算,理論上能讓重複場景的等待時間大幅縮短。
- 記憶體階層管理:這台機器記憶體特別大,規劃讓常用的資料留在高速區域,不常用的分支往外挪到系統記憶體甚至硬碟,讓長文本、多工作階段可以撐得比一般消費卡更久。
- 執行底座(核心挑戰):既然已經找到「打雜指令太多、張量核心沒吃飽」這個病灶,那就設計一套專用的運算格式——在模型載入的時候,先把量化過的權重資料在 CPU 上重新排列好、攤平,讓格式直接對齊硬體讀取的方式,執行的時候不用再即時去拆解、換算,直接把資料灌進張量運算單元。當時抓到一個細節:那個「拆解、換算」的動作,理論上重排後幾乎不用花時間(實測 repack 吞吐約 277 GB/s,佔用時間微乎其微)——照這個邏輯算,把這一大塊打雜指令消滅掉,矩陣運算的效率應該能從原本的水準往上跳一大截。
這套規劃聽起來很完整,野心也很大,花了一整天編譯測試、編譯測試數十次後的結果——
現實給的第一巴掌:野心最大的那塊,反而摔得最重
三層規劃裡,前兩層(前綴快取、記憶體管理)算是設計文件層級的東西,還沒真正投入大量時間去實作;真正花時間動手做的,是第三層(執行底座)那個「重新排列量化格式,直灌張量核心」的專用運算核心。
這是整個自研計畫裡最有信心的一塊——理論算得漂亮,格式重排的成本理論上趨近於零,消滅打雜指令換來的效能提升,理論上該是最大的一筆。
結果做出來上機一測:ffn_up 從 3947 微秒變成 4195(慢 6%),ffn_out 從 4142 變 4360(慢 5%)——比什麼都不改的原始版本還慢。
這一下摔得挺重的,原本以為找到的病根、算出來的帳,竟然完全兌現不了。回頭拆開來查才發現三個問題疊在一起:
- 省下來的那些指令,其實原本就藏在記憶體等待的空檔裡,根本沒有真的拖慢速度,砍掉它們自然也不會有感覺;
- 重新排列後的資料格式反而變胖了——每行 288 位元組 vs 原來 176,胖了 64%,多花的記憶體頻寬,比省下來的運算時間還要多,得不償失;
- 量化格式裡有一段數學運算,本質上沒辦法簡化——原本以為 scale 可以提出來一次算好、重複使用(Post-scaling),但仔細算過才發現兩個 scale 都隨循環變數變化,沒有公因子可提,是個死巷子,每個元素都得個別算,沒有捷徑。
這才是這兩天最挫折的一個時刻——不是一個小失誤,是整個計畫裡最被寄予厚望、算得最有把握的那一塊,徹底摔了個大跤。
接下來一連串,野心勃勃的三個計畫,一個一個被現實打臉
除了執行底座這個大摔跤,接下來嘗試的其他方向,也幾乎每一項都被推翻:
- 懷疑排程效率是主因——開同步模式(
HIP_LAUNCH_BLOCKING=1)重測,pp8K 跑出 332.5,反而比正常的 317.8 更快,證明 launch overhead 根本不是瓶頸,真正的瓶頸在矩陣運算本身的效率太低;
- 拿現成調校工具(hipBLASLt)想要優化矩陣運算,單算子從 21.0 量到 33.4 TFLOPS(1.59x),理論上該起飛——結果上機實踐(Q5_K repack)反慢 5~6%,端到端 pp512 也掉了約 2%。因為比較方式漏算了隱藏開銷:純粹局部執行時量測出數字沒含格式轉換與多出來的頻寬帳,端到端一算,1.59x 兌現不了;
- 想省掉重複讀取的浪費——GDN staged kernel 寫完、正確性沒問題,但收益從沒兌現(env 門都沒開);KV 轉 q8 測完 320.2 vs 317.8,純誤差範圍,直接關閉。直覺上會認為 GPU 核心反覆向顯存請求讀取了同一個變數高達 32 遍,造成了嚴重的頻寬浪費;但實際去寫 Kernel 想省掉它時,才發現這段浪費早就被硬體快取接住了;
- 想解開 MMQ occupancy 上限、fattn tile 這類看似隨手的參數限制——一放開就跑出錯誤數值,後來才發現這些限制底下是沒修好的硬約束(fattn 的 head_dim=256 上不了 tensor 路徑、MMQ 的暫存器形狀),是道承重牆,不能亂動,只好原樣焊回去;
- **既然 WMMA256 跑不動、但現成的 WMMA-128 跑得很順,那我能不能把維度為 256 的向量直接「一刀切成兩半(前半部 128、後半部 128)」?分開丟進兩個 WMMA-128 算,算完再拼起來,不就繞過硬體限制了?但Muse Spark提醒我,把兩半個別算完 Softmax 的結果就算拼回去或加起來,在數學上根本不等於原本「完整的 A+B 算出來的 Softmax,及時懸崖勒馬,零成本撤回,沒有浪費工夫去真的實作。
最後,靠著幾個零碎的參數微調(MMQ occupancy 候選 + tile nbk64 + 關 IOMMU),兩天下來總共拿到了 10.4% 的實測效能提升(pp8K 317.8→345.2,pp32K 263.5→290.9),把差距從 2.17 倍收斂到 1.96 倍。不是無效,但離「自己做出一個能打贏對手的引擎」還遠不夠看。
轉念:與其自己硬幹,不如先看看對手做了什麼
眼看自己這條路已經摸到天花板,我連27B的挑戰都搞不過,更何況更難的125B-Flash-Next,乾脆換個角度——把 halogen 的執行用SASS拆開,看看裡面組合語言層級的指令成分長什麼樣子,這中間又讓AI Agent持續執行、解包、數字解析、翻譯理解等循環又花了我一天的時間。
手法:podman pull 拉容器 → 匯出檔案系統 → 引擎 binary 裡的 .hip_fatbin 段 → 切出 HSACO → 反彙編。0.1.3 版 8 個 code object,0.6.1 版 3 個、1047 個 kernel,全攤開。
要先講清楚一件事:這不是「反組譯出原始碼」,更不是破解或竊取——這只是看執行檔被編譯成什麼樣的底層機器指令,跟看一棟房子的磚頭怎麼疊沒有本質差別,看不到設計圖、看不到背後的邏輯,只能看到最後蓋出來的成品長什麼樣。
看完之後,有幾個觀察到的現象:
觀察到的:
- halogen 的核心運算(
k_gemm_i4,熱區 298 條指令裡 128 條是 WMMA、佔比 43%),張量運算指令跟打雜指令是 0.5:1,官方路徑是 5:1——打雜佔比低一個量級;
- 它用的是一條查不到公開文件的整數矩陣運算指令
v_wmma_i32_16x16x16_iu4——注意是 iu4 不是 iu8,單條做 8192 個乘加,是官方那條 iu8 的兩倍。攤到每個乘加操作,halogen 的打雜只有官方的 1/43;
- 暫存器用量 ~238(VGPR)+ ~48(SGPR),遠超一般教科書建議的安全範圍(32),這通常會讓 GPU 沒辦法同時排很多工作,但看起來作者是刻意這樣做——熱區 128 條 WMMA 深度展開,換取單一工作內部連續執行更多指令,用超長直線程式碼蓋住記憶體延遲,不用頻繁切換等待;
- 從記憶體存取模式判斷,0.6.1 世代的手寫 linear 全是純 ALU(給 decode/小矩陣),prefill 主力是解包 kernel + 調優庫兩遍式——權重資料可能在載入時就先重新排列過,執行時不需要再做即時的格式轉換(0.7.0 起官方也親口說 BYO-GGUF 會無損重排,可互相印證);
以下AI Agent順便去爬文作者的ChangeLog:
- Flash-Next 專用系統級調度(0.6.3版):MoE 計數排序與 NVMe 並行讀表:針對 Qwen3.8-Flash-Next 版本,對手進一步實作了開源框架完全不具備的系統級優化,Qwen 的 47.7 GiB 查N-Gram找表不進 RAM而擺SSD,改用 Linux page cache。以 CPU開64 線程平行聚合讀取,將長 Prompt 讀表時間從單執行緒的 50 秒壓制到 1.3 秒 。
- 一整套Decode加速專用的Drafter(0.5版):作者poenist.ai自己搓了一個 MTP head fine-tune、另一個 DFlash2,block drafter,5 層 ~2.2B 參數,讓MOE架構下也能讓接受率達到97%以上。
- 注意力層面的動態微區塊剪枝(0.1版):手刻 QSA Micro-block 稀疏Qwen3.8 原生架構是「3:1 混合注意力」(75% 線性 GDN + 25% 全量 Attention)。開源框架在面對那 25% 的全量層時硬算 O(N^2),在 128K context 下 Attention 吃掉 65% 時間,從指令模式判斷(
k_attn_sel:613 條、零 WMMA、26 個數據依賴分支,輸出索引;k_attn_qs_bt* 幾十個變體吃 int* block 索引;k_attn_fa2 模板參數就叫 FaSel),它針對長文本的注意力計算做了動態 block 篩選,跳過不重要的部分,而不是每個 token 都跟所有歷史硬算一遍——這也能解釋為什麼差距隨深度拉大(128K 差 2.78x):GEMM 調優只能給全段固定倍率,越拉越大的那截全是 attention,但這QSA推論引擎可靠嗎?
推論引擎級的QSA精度有沒有偷工?這部分我真的自己動手測了。
如果 halogen 真的靠「跳過部分計算」來加速,會不會犧牲了準確度?這個不能靠猜,得自己動手驗證。
我寫了幾個測試腳本,直接對它的線上服務做壓力測試(每次填充文隨機生成,防它的 prompt cache 作弊):
- 在很長的文本裡(16K/32K/64K/128K 四檔 × 深度 5%~95%),埋入隨機的識別碼,測試模型能不能準確找回來——52 次測試,52 次全部命中(含 128K@50% 死角 3/3);
- 設計一連串需要跨段落記憶的變數運算題目(A 在 15%、B=A^mask 在 52%、尾端問 B),測試模型能不能正確追蹤——一開始 成功率0/4,但Agent仔細檢查只是單 nibble 心算翻車;引導它把 A、mask、B 分步講出來後答案完全正確,證明關鍵資訊沒有被漏算,是算術失手不是撿不到;
- 128K 長度的程式碼考題(424K 字元、約 121K tokens,六個寄存器值埋 5%~90%,每段一個真值一個誘餌),埋入需要精確記住的變數數值——6/6 全對、零誘餌,160 秒交卷。
**結論:至少在這些測試裡,沒有看到精度被犧牲的跡象。**這部分是我自己動手驗證出來的第一手結果,不是猜測,可信度最高。(順帶一提:halogen 官方自報的 NIAH 是 148/150 = 98.7%,跟我的 52/52 互相印證。)
看完這些觀察報告後,我直接投降.....自己即使以AI-Agent搭配C++逆向工程試搓並測試完初版起碼也要幾天還不一定驗得出效果,等到驗證完後Halogen也開源了,別自己瞎搞....
這兩天,到底值不值得
老實說,我沒有做出一個打贏 halogen 的引擎,10.4% 的提升那只能說是調校...跟對方整盤翻起來重寫的優勢比起來只是零頭。但這兩天讓我親手驗證了幾件事:
- 這顆晶片(RDNA3.5)在長文本推論這種需要大量資料搬運跟精細排程的任務上,確實有一些先天的架構限制,不是靠調參數就能繞過的——MMQ 的 21 TFLOPS 天花板、62% 承重牆打雜、81 VGPR 的形狀,說的都是同一件事。這也印證了開頭提到的大背景:RDNA 這個世代在張量運算規格、資料搬運機制上,跟同世代 NVIDIA 硬體本來就有結構性差距,兩天的調校頂多是在這個先天限制底下,把軟體沒做好的那部分補一點回來(+10.4%),補不了硬體規格本身的落差;
- 靠 AI 工具輔助,一個人也能在兩天內走完「規劃→動手→失敗→驗屍→重新規劃」的完整循環,而且每一次的失敗都留下了明確的原因(-6% 的驗屍報告、偽標準的撤告、數學死巷子的證明),不是白費工夫;
- 然而Halogen雖然規模小(peonist-ai 個人作者,GitHub 396 star,引擎閉源),但確實摸到了超多AMD官方路徑沒做到的細節優化:iu4、238 VGPR 換 ILP、稀疏注意力整套、特製Drafter、NVMe多線查表——沒有一條是一點訣的魔法,全是工程;
- 作者在社群裡也提到,等這套系統的架構整理完,會考慮全面開源。如果那天真的到來,這兩天累積下來的這些觀察跟教訓,應該能幫上想深入了解RDNA3/4家族的人一些忙。
以上內容基於個人實測與觀察所做之推導,對方原始碼尚未開源,僅供參考討論。