GPU Mode 與 Core Automation 合作,最近舉辦了一場以自動研究為主題的競賽。問題是實現批次方形緊湊 Householder QR 分解,也就是 QR 分解。我在 183 名參賽者中排名第 12 名,最終比基準解決方案快了 232 倍。這篇文章將介紹我是如何達成這個目標的。我將會說明我的方法、學到的東西以及在競賽中遇到的瓶頸。這是我第一次認真嘗試自動研究。有些人會稱之為「迴圈工程」,坦白說,這樣稱呼也沒關係。
請注意,您不需要詳細了解數學或問題本身,就能理解這篇部落格文章的大部分內容。我著重於我的方法,而將數學和問題本身放在次要位置,因為大多數讀者可能沒有參加過這場競賽。
您可以在此查看完整的競賽頁面:問題連結與排行榜
這場競賽是 GPU Mode 的「研究時代的線性代數核心」系列的一部分。
我們收到一批批的方形 FP32 CUDA 矩陣 A,形狀為 batch x n x n,目標是返回與 torch.geqrf(A) 相同的緊湊 Householder QR 表示法:一個 H 矩陣,其上三角形是 R,下三角形儲存 Householder 向量,外加一個反射係數的 tau 向量。檢查器使用 torch.linalg.householder_product(H, tau) 重建 Q,將 R = triu(H),並驗證:
A ≈ Q R, QᵀQ ≈ I, QᵀA ≈ R
在正確的提交中,排行榜根據各種形狀和條件案例的幾何平均運行時間進行排名。重要的尺寸是批次方形矩陣,例如 512 x 512,還有更大的 1024、2048 和 4096 的案例。允許內部使用低位元 FP16、FP8 或 NVFP4,但返回的因子仍必須滿足 FP32 風格的 QR 檢查。
這裡 Q 是正交的,意味著它的列向量長度為單位向量且彼此垂直,而 R 是上三角形矩陣,意味著對角線以下的元素都為零。競賽並沒有要求我們直接輸出稠密的 Q 和 R;它要求的是緊湊的 Householder 表示法,讓檢查器能夠重建 Q 並從上三角形讀取 R。
對於上面 3x3 的例子,第一個反射器一次就將第一列 (12, 6, -4) 映射到 (-14, 0, 0)。-14 變成了 R₁₁。其工作原理在數學部分有詳細說明。
GPU Mode 提供 popcorn CLI 工具給參賽者,使其易於使用。參賽者可以使用此工具直接測試、基準測試並提交到排行榜。檢查器也提供了按形狀回饋以及整體幾何平均時間。
敏銳的觀察者會注意到,這是一個非常適合編寫迴圈的設置。參賽者渴望緊密的迴圈回饋。這讓他們可以隨心所欲地進行爬坡優化。
GPU Mode 的競賽通常會提供一些方法來迭代核心。您可以直接提交,或者像 Modal 這樣的贊助商會提供額外的點數。在這裡,組織者基本上允許無限次提交,只要您錯開時間。如果您不這樣做,隊列就會變長,所有人的運行都會超時。一度,工作空間甚至用完了 Modal 的點數,因為每個人都在瘋狂提交。這是一種讓學習更容易的好方法。
在 14 天的過程中,我提交了超過 1500 次。
我對 GPU 核心優化的基礎知識(主要在 Triton 中,並對 CUDA 有一些了解)已經有一年了,但沒有在這方面從事過專業工作。我想告訴您的是,在排行榜上,我身邊的人都比我強。我上面一位的參賽者(CUDA Colonel)是 NVIDIA 的首席工程師。
撇開光環效應不談,由於我了解基礎知識,並且最近閱讀了 GatedDeltaNet 的文章,所以我對一般的 GPU 核心術語比較熟悉。
您對某件事了解得越深入,就越能更好地提示大型語言模型,因為您可以將未知的未知轉化為已知的未知。
同時,值得注意的是,即使沒有領域知識,這個競賽也是可以完成的——您可能無法進入前 10 名,但僅僅依靠您的框架/代理迴圈或其他方式,就可以獲得比基準線可觀的加速。
我在競賽中的第一步是了解什麼是 QR 分解以及如何進行。有很多方法可以做到,例如 Gram-Schmidt 和 Householder 反射。競賽強制要求使用 Householder 反射。我與 Claude 來回討論,並觀看了幾個 YouTube 影片來建立直覺。在與 Claude 討論後,很明顯我們需要使用區塊化 Householder 演算法作為主要架構,並配合後續的 WY 更新。事實證明,GPT-5.5 對此也有一個不錯的想法。QR 分解是一個相當知名的問題。
我發現這個概念很有趣,因為矩陣分解出現在許多現代 LLM 訓練優化器變體中,特別是在使用矩陣預處理的方法中,例如 Shampoo 風格的優化器和相關方法。Muon(Kimi 使用)是另一個很好的例子:它不是將權重更新視為一個巨大的展平向量,而是保留矩陣結構,並對動量更新進行正交化,通常通過幾次 Newton-Schulz 迭代來近似極分解。
如果您對數學感興趣,我建議您快速瀏覽本節,否則可以隨意跳過。唯一需要注意的是,Householder QR 中存在順序依賴性,這使得 GEMM 變得困難。我們使用區塊化 Householder 來使其更接近矩陣乘法的形狀。
快速回顧合約:輸入是一批方形 FP32 矩陣 A;輸出是 torch.geqrf 返回的緊湊 (H, tau) 格式。H 的上三角形是 R。對角線下方,H 儲存 Householder 向量,tau 儲存每個向量的單個標量。檢查器從 (H, tau) 重建 Q 並驗證 A ≈ QR。
暫時忘記矩陣。在浴室鏡子裡,您的倒影與您在鏡子前的距離一樣遠,是直線的。
如果 x⊥ 是 x 垂直於鏡面的部分,反射就是將該部分減去兩次:
x' = x - 2 x⊥
所以 Householder 反射就是找到垂直部分並將其減去兩次。
Householder 向量就是鏡子,以緊湊的方式儲存。在程式碼中,我們不傳遞整個鏡面。我們儲存一個向量 v,它垂直於鏡面。鏡面是所有垂直於 v 的部分,而反射沿著 v 移動。
垂直部分只是 x 沿著 v 的投影,即 vᵀx / vᵀv 份的 v。將其代入上面的減法:
Hx = x - τ v (vᵀx), τ = 2 / (vᵀv)
所以 tau 就是 2 / (vᵀv):2 這個因子和 v 的長度打包成一個預先計算好的數字。v 選擇鏡面,tau 縮放更新。(我將數學反射器寫為 Hⱼ,並保留 H 作為緊湊輸出矩陣。)
QR 的目標是將 A 轉換為上三角形矩陣 R。第一列應該變成類似 (*, 0, 0) 的形式,第二列應該在第二行以下為零,依此類推。
Householder 鏡面很有用,因為它可以一次性處理一列。以上面 3x3 的例子中的第一列為例:(12, 6, -4)。我們想將它變成 x 軸,使下面的元素為零。反射只能改變方向,不能改變長度,所以目標的長度也必須是 14。一個有效目標是 (-14, 0, 0)。在進行該反射後,6 和 -4 的元素就消失了,這正是我們想要的。
我們如何找到鏡面?它位於該列和目標之間的中點,所以 v,即穿過鏡面的向量,就是該列減去其目標:
v = (12, 6, -4) - (-14, 0, 0) = (26, 6, -4)
然後計算 tau = 2 / (vᵀv)。反射器本身是:
H = I - τ v vᵀ, τ = 2 / (vᵀv)
Hx = (I - τ v vᵀ) x
Hx = x - τ v (vᵀx)
H A_active = (I - τ v vᵀ) A_active
H A_active = A_active - τ v (vᵀA_active)
這將當前列變為 (-14, 0, 0),並一致地重寫其他列,以便下一個反射器從更新後的矩陣構建。
我們對每一列重複執行此操作。粗略表示,重複的更新如下所示:
A⁽¹⁾ = A⁽⁰⁾ - τ₁ v₁ (v₁ᵀA⁽⁰⁾)
A⁽²⁾ = A⁽¹⁾ - τ₂ v₂ (v₂ᵀA⁽¹⁾)
A⁽³⁾ = A⁽²⁾ - τ₃ v₃ (v₃ᵀA⁽²⁾)
每一行都使用前一行產生的矩陣。每個反射器都會將其列對角線以下的元素歸零,而不會干擾已經完成的列。最後一個反射器之後,A 就變成了上三角形的 R:
鏡面不會改變長度或角度,因此每個 Hⱼ 都是正交的,它們的乘積 Q 也是正交的。這就是檢查器驗證的正交性免費獲得的來源。
一旦處理完第 j 列,其對角線以下的元素就成為死空間。geqrf 重用這些空間來儲存 vⱼ 的尾部(第一個 1 是隱含的)。對角線及其上方是 R;對角線下方是反射器;tau 作為單獨的向量隨行。這就是為什麼檢查器需要 H 和 tau 來重建 Q。
Householder QR 一次一列地將 A 的對角線以下元素歸零。每一步都從當前列構建一個反射器,並將其應用於右側的所有元素。問題在於反射器 j+1 是從反射器 j 已經作用過的矩陣構建的。所以您不能重新排序步驟,也不能合併它們。這是串列的,而串列的矩陣向量工作在 SM 的慢速向量通道中運行,而張量核心則閒置。
經典的解決方案是區塊化演算法。您選擇一個由 b 列組成的狹窄面板(例如 32 或 64 列),並在其中執行所有串列工作。這沒關係,因為面板只有 b 列寬,所以成本仍然很低。然後,不是一次將面板的 b 個反射器應用於矩陣的其餘部分,而是將它們壓縮成一個單一的秩-b 更新(「WY 表示法」),並一次性用三個連續的矩陣乘法處理整個尾部區塊。串列工作仍然限制在面板內,而其他所有內容都變成了 GEMM,這正是張量核心想要的形狀。
具體來說,WY 表示法將面板的 b 個反射器壓縮成一個單一的秩-b 更新。將面板的 Householder 向量堆疊成 V 的列向量 [ v₁ , v₂ , … , v<0xE2><0x82><0x99> ],構建一個小的 b x b 上三角形矩陣 T,然後:
WY 表示法:H = [ I - τ v₁ w₁ᵀ | -τ v₂ w₂ᵀ | … | -τ v<0xE2><0x82><0x99> w<0xE2><0x82><0x99>ᵀ ]
尾部區塊更新變成三個 GEMM 形狀的步驟:
Z = T Vᵀ
Q_trail = I - V Zᵀ
A_trail ← A_trail Q_trail
面板是串列矩陣向量工作被卡住的地方,但它只有 b 列寬。它右側的所有內容都是大的尾部區塊,這是純粹的 GEMM。隨著面板沿對角線移動,尾部區塊會縮小。
如果您想理解數學,我建議您與 Claude 一起腦力激盪。也可以看看 Mike 的文章,他在其中以比我更具描述性和視覺化的方式涵蓋了數學。他在競賽中獲得第 5 名,並分享了他對問題的見解。
另外兩件事被證明是具有挑戰性的:可靠地在內部使用低精度(尤其是在處理病態輸入時),以及應對各種形狀(n = 32、176、352、512、1024、2048、4096)和批次大小的廣泛分佈,其中最大的矩陣批次太少而無法填滿張量核心,而 n = 32 又太小,我們必須將許多矩陣打包到一次核心啟動中。
在獲得基本理解後,我要求 Codex 進行基本設置:添加 problem_statement.md,在 AGENTS.md 中提及提交和使用 popcorn CLI 的基本細節,並維護一個 log.md 來記錄提交及其狀態(接受/拒絕以及按形狀的時間)。如果與 Karpathy 的自動研究進行對比,我的 AGENTS.md 和 problem_statement.md 最初就是我的 program.md。
日誌作為有效和無效想法的證據。未來的代理會話可以閱讀日誌並快速檢查是否嘗試過某個想法。在 3000 µs 之後,我投入了更多精力記錄,因為事情開始變得更難。
Codex 的酷之處在於,您可以告訴它做什麼,它就會實際去做。如果您給它一個足夠詳細的目標提示,它可以工作數小時。它知道所有的數學和程式碼,並且擁有它所需的所有回饋。我有這個直覺,但我仍然驚訝於 GPT-5.5 能將效能推升到超出基準解決方案(即 torch.geqrf/cuSolver 函數)的程度。
我最初的幾次會話是手動提示,要求在 Triton 中實現解決方案並針對 n = 512 和 n = 1024 的形狀進行優化,因為它們在最終的幾何平均值中佔有最大的權重。
但是,如果您想讓模型一直迴圈直到達到某個目標,請使用 /goal。您可以給出具體、可實現、量化的目標。我發現給出一個好的數值目標,然後跟隨具體的標準,效果很好。
範例:「僅使用 Triton 或 CUDA,並打破我們目前為止 n = 512 的最佳時間。嘗試幾種想法,方法是直接提交到排行榜或使用 Modal 分析。在新的一組實驗中完全移除 cuSolver。我們只將其作為備用方案。」有一次,這個目標運行了一天多。
我每 2-3 小時輸入一次,以引導模型朝我想要的方向發展。有些日子,我讓它在沒有監督的情況下過夜運行。最初幾天,我的指示主要圍繞著引導模型嘗試不同形狀的不同想法(稍後詳述),要求它使用更多 Triton,更少 PyTorch,以及更少的備用方案!我學到的一課是:我經常想頻繁檢查我的代理,但您必須信任它並讓它完成工作。您必須讓開代理的路;只有當它卡住時,才將其引導回來。
在使用 /goal 時,您可以通過使用 /btw 或 /side 來提問,而不會中斷迴圈。這會創建一個臨時線程,其中包含您主對話的上下文。我喜歡這種檢查代理並提供監督的方式。
我會問這樣的問題:兒子,你贏了嗎?當前活躍提交中的演算法/變更是什麼?向我解釋這個概念。您目前正在處理哪些想法?進展如何?最近有什麼突破?我們可以將其轉移到其他形狀嗎?您接下來的最佳想法是什麼?根據回應,我會諮詢 Claude 以增進對相關概念和瓶頸的理解,然後將我的想法提供給 Codex。提問後,我會回到主線程並傾倒想法來引導模型。
如果 /goal 或某種類型的迴圈正在運行,那麼您可以將指令排隊到 codex 中提問,或者執行 /btw。(圖片是我在檢查 codex,詢問它在做什麼,下一個想法是什麼等等)。如果您按下 esc 中斷迴圈,則 /goal 會暫停 https://t.co/Yx6EpH5PKo pic.twitter.com/imqZc6RPkz
基準的 torch.geqrf 路徑總體約為 419 毫秒(419,000 微秒)。在實現了該形狀的區塊化 Householder 路徑後,我一天內就達到了 n = 512 形狀的 5000 微秒。這是 geomean 中權重最大的形狀。
上面的 232 倍數字來自比較粗略的 419,000 微秒基準線與最終追蹤到的 1,805 微秒結果。下面的演進圖從我追蹤的提交歷史記錄中恢復的第一個點開始,因此它顯示了後來的 108,803 到 1,805 微秒的弧線,而不是完整的基準線到最終的比例。
優化核心時,讓工作更具矩陣形狀,使張量核心停止閒置,是您的生命目標。
在 3000 微秒之後,優化變得更加困難。我必須更多地參與到迴圈中,學習概念並引導模型。我授予 Codex Modal 分析權限,讓它運行 torch 分析 / nsys 分析來測試不同的想法,比較實現,並更快地掃描參數。(後來,組織者也提供了一種進行 NCU 分析的方法。)
QR 核心的粗略結構演進,從基準/函式庫密集的路徑,到自訂面板工作、融合組件和 GEMM 形狀的尾部更新:
在分析整個形狀並確定瓶頸方面,來回進行了很多工作。對於這個問題,啟動開銷和面板開銷佔主導地位。我們很少受記憶體或計算限制。大部分精力都花在尋找減少面板開銷並使 WY 更新更易於 GEMM 的方法上。
我發現自己不斷問的問題:
什麼是面板開銷/啟動開銷?我們如何解決它?
我們缺少哪些資訊,如何獲取?
最新的分析結果是什麼?您有什麼看法?
尋找融合候選。是否有 Triton 融合級別的候選?
尋找可能的編譯器優化。我們可以將哪些運行時信號轉換為靜態信號?想法是給編譯器提示。我曾因這個原因獲得了 200 微秒的提升。
要求它幾次查看 Triton 生成的編譯成品。
有哪些數值技巧可以利用精度餘裕?
您能否部署子代理來進行一些數學運算並找到可能的優化?
部署代理來搜尋可能拖累我們的錯誤。
是否有可能的歸約融合?我看到 Codex 發現了一個,然後我經常開始猛攻它。歸約是進行求和或尋找最大值等操作。由於需要迭代整個序列來完成工作,我們可以一次執行多個這樣的操作。
在 3000 -> 1800 微秒的範圍內,我開始面臨的一個主要挑戰是模型陷入局部最大值。這看起來像是無休止的手動調整參數和相同想法的小變體。
幾年前,代理會因為不夠聰明或知識不足(或驗證迴圈不夠健壯)而陷入(末日)迴圈。人們過去會嘗試採樣、溫度變化和不同的解碼策略來解決這個問題。
然後隨著時間的推移,模型變得更聰明,預訓練了更新的知識,我們獲得了大量的 RL 擴展和推理時間計算(訓練模型以更多的 token 進行思考)。
現在模型在尋找新想法和「研究品味」方面遇到困難——考慮到驗證器回饋、先前證據和我們的評估(在我們的案例中是分析器回饋),下一個最佳想法/實驗應該是什麼?好的想法生成是下一個偉大的冒險。我最近也寫了一篇關於這個主題的文章。
我使用了以下策略來幫助模型擺脫局部最大值:
候選者光束:很長一段時間,我做了一件愚蠢的事。我只保留一個最佳候選者,然後用新的候選者來測試。如果代理嘗試一個新的結構性想法或一個顯著的重大變更,那麼它得分低於我們當前最佳提交的可能性很高。然而,經過幾次迭代,該變更可能會優於我們的最佳候選者。基於這個觀察,我引入了一些維護 3-5 個候選者光束的指令。
人類介入迴圈:當模型長時間卡住沒有進展時,我會充當秘密醬料來引導模型。
鼓勵模型冒更大的風險並嘗試雄心勃勃的想法。您可能不相信,但這奏效了。此外,Claude 在幾輪之後經常放棄,並以「我們已經用盡了所有優化來達到 x geomean」等藉口;另一方面,Codex 則更堅持。
使用更強的顧問模型來產生更多樣化的想法。在我的框架中,這看起來像是 AGENTS.md 中的一個指令,鼓勵模型使用無頭調用 claude -p 來獲取想法、提供分析數據等。
指示模型經常使用子代理來嘗試雄心勃勃的想法,搜尋網路上的部落格和論文,瀏覽我上面列出的列表,並尋找微優化。
使用 NCU 和 Modal 進行分析,比較結果,並處理 Modal 和 NCU 分析都指出的瓶頸。
我曾考慮過一種多代理群體方法,但沒有嘗試。
我認為強顧問策略(如 GPT-5.6 Sol 或 Fable)將成為自動研究流程中的標準策略。考慮非常大的模型和最先進的訓練。Claude Code 也為此提供了 /advisor 命令。
我們將顧問策略引入 Claude 平台。將 Opus 作為顧問,與 Sonnet 或 Haiku 作為執行者配對,以更低的成本獲得接近 Opus 級別的智慧代理。pic.twitter.com/fRkegyMs5t
到最後,我的目錄結構看起來像這樣:
到最後,我的 AGENTS.md 看起來像這樣:
此工作空間用於排行榜風格的優化工作。除非使用者明確覆蓋,否則請遵循這些固定指令。
請勿進行單一佔優的爬坡優化。維護一個活躍想法家族的小光束,以便局部負面結果不會過早地扼殺有用的組件。
雖然我的簡單框架有改進空間,但我的大部分缺點都是特定於問題的。在掃描了前 10 名提交並閱讀了 Mike 的文章後,我有了這些體會。
n = 512 和 n = 1024 的案例有多種輸入分佈,例如稠密、聚集、秩虧損、混合和近秩族。更快的核心已經編寫了數據檢測器並利用了分佈,例如低秩案例有很多零,那麼我們如何利用這一點呢?我本可以在這方面進一步推動 Codex,或者至少多詢問這方面的事情。
前 10 名的解決方案更積極地移除了函式庫函式。例如,第 2 名和第 5 名的解決方案使用了自訂三角反矩陣,而不是使用 PyTorch 的三角求解器。我的解決方案在 PyTorch 和 Triton 之間來回進行了很多操作。
本來可以將尾部矩陣保留在 fp16 中,而不是在不同表示法之間反覆移動。這對我來說是一個未知的未知,純粹是領域專業知識的缺失。
我一開始就應該有候選者光束。
編寫更健壯的分析器來測試混合精度案例。
未能使用 tcgen05 指令,即 NVIDIA 的第五代張量核心指令,在 Blackwell 上利用 B200 張量核心。
我在 183 名參賽者中排名第 12 名,比基準線快了 232 倍。更重要的是,我學到了很多關於 GPU 核心優化、自動研究(或迴圈工程?)基礎知識以及一些 B200 特定細節。
我也得出結論,領域專業知識可以加速框架設計和人類介入迴圈的引導。我還有一些觀察。工程師希望將事物通用化的程度存在一個光譜。在左側,人們希望創建一個可以自我改進的通用框架。在另一端,人們對創建非常特定於問題/環境的框架感興趣。我個人傾向於特定於問題的框架。
希望您喜歡閱讀這篇文章並學到新東西。如果您喜歡這篇部落格,請點讚/投贊成票/分享!
對了,該系列的第二場競賽,特徵值分解,目前正在進行中。排行榜上見。
我使用了以下資源來修改或學習新概念。
如何優化 CUDA Matmul 核心以獲得類似 cuBLAS 的效能:工作日誌
在 H100 上超越 cuBLAS:工作日誌
在 NVIDIA B200 GPU 上進行快速 QR 分解 - Mike 的文章幫助我修改了問題本身 lol。
額外希望深入研究的參考資料:
Simon 的部落格 - CuteDSL 和 B200 特定內容。
gau-nernst 的部落格 - gau-nernst 獲得第 2 名。
Rohan Anil、Mark Saroufim,感謝您舉辦競賽。Sinatras 感謝您偵測到獎勵漏洞並給予我最初的鼓勵。
Mike 獲得第 5 名,並撰寫了一篇精彩的文章,解碼了問題並包含很棒的視覺化效果。
Levidiamode 感謝您提供有關 GPU 核心的寶貴文章,Tokenbender 感謝您鼓勵我寫這篇文章,CUDA colonel 感謝您在 GPU Mode 伺服器上提供提示。
Claude Opus 4.8 和 GPT-5.5 (Codex) 在編輯和數學相關寫作方面提供了幫助。