GPU 工具鏈的效能比較,長期偏愛矩陣乘法、卷積與歸約等規則運算。這類工作可預先知道每個執行緒要處理多少資料,記憶體也容易切成整齊區塊;然而機器人建圖、稀疏資料結構與動態工作佇列面對的,往往是另一個世界。
2026 年 8 月提交至 arXiv 的單一作者預印本〈What Irregularity Costs〉,以同一套雜湊分塊 TSDF 融合核心,比較手寫 CUDA C++、透過 NVIDIA cuda-oxide 編譯的 Rust,以及 Triton。這是尚未經同儕審查的實驗結果,不應讀成通用語言排行榜;它真正有價值之處,是指出工具抽象一旦遇到不規則控制流,成本可能從百分之幾突然放大到數十倍。原始論文:https://arxiv.org/abs/2608.08287
TSDF 為何會把工具鏈逼到牆角
TSDF(截斷式帶符號距離函數)把場景切成體素,記錄每個體素至表面的帶符號距離;表面位於距離跨越零的位置。若把整個三維空間密集配置,記憶體需求很快失控,因此系統只配置感測到表面的區塊,再以空間雜湊索引。論文中的區塊含 8×8×8、共 512 個體素,雜湊表採開放定址與線性探測。
融合被拆成兩階段。allocate 階段要沿深度點的視線穿越截斷帶、計算區塊座標、探測雜湊表,必要時再以 compare-and-swap(CAS)搶占空槽;每條執行緒探測幾次並不相同,還可能爭用同一位置。update 階段則多半第一次查找就命中,接著累加距離、權重與色彩。前者是不規則控制流,後者接近 GPU 工具最熟悉的規則運算。
結果不是「誰最快」,而是成本在哪裡裂開
來源事實:論文在兩張 Blackwell 架構顯示卡、19 組工作負載上,固定資料布局、雜湊函式與 256 執行緒啟動配置。相較 CUDA C++,Rust 在不規則 allocate 階段為 1.02 至 3.34 倍,Triton 為 11.2 至 31.6 倍;規則 update 階段則分別為 0.96 至 1.19 倍與 1.1 至 2.6 倍。若只看完整融合路徑,allocate 約占十分之一,Rust 在論文所稱的真實深度資料上僅慢 1% 至 3%,Triton則約為 2.6 至 2.9 倍。
合理推論:產品端只看端到端平均值,容易掩蓋某個低占比核心的結構性問題。當資料量、碰撞率或處理管線比例改變,原本只占十分之一的瓶頸可能浮上來。因此比「整體慢幾倍」更重要的問題,是效能差距是否隨爭用程度擴張。
Triton 的代價來自不能提前退出
CUDA C++ 與 Rust 可讓已找到鍵值或空槽的執行緒離開探測迴圈;論文中的 Triton 實作卻必須執行到編譯期設定的固定上限,再以遮罩標記已完成的 lane。更棘手的是,Triton 官方文件列出的 atomic_cas(pointer, cmp, val, sem, scope) 參數中沒有 mask,已完成的 lane 仍需被導向額外的暫存區,避免誤寫正式雜湊表。這項介面事實可由官方 API 交叉驗證:https://triton-lang.org/main/python-api/generated/triton.language.atomic_cas.html
這不只是速度問題。固定探測上限若設得太小,碰到較長探測鏈就會漏掉區塊;設得太大,又會增加編譯時間、二進位大小與無效工作。論文記錄在雜湊負載因子 0.283 時,Triton 建出 1,147 個區塊,而 CUDA C++ 為 1,160 個,且有 27,943 次貢獻被丟棄。這類失敗沒有錯誤訊息,最後表現為重建表面缺洞,比單純延遲更危險。
Rust 接近 CUDA,卻揭露安全與快取的張力
Rust 的差距不是較多指令或較低占用率。論文的硬體計數器反而顯示,Rust 使用較少暫存器、分支與 CAS,卻有較低的 L1 命中率及更多 L2 流量。作者追查後認為,原因是以 GPU scope 原子載入讀取其他串流多處理器可能更新的位置;為維持跨 SM 可見性,這類讀取不能依賴不具跨 SM 一致性的 L1。
NVIDIA 官方記憶體模型確認,device thread scope 涵蓋同一 GPU、同一記憶體同步網域內的 GPU 執行緒,支持論文對同步範圍的基本描述:https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cuda-cpp-memory-model.html。但「因此繞過 L1 並造成多少損失」仍是該論文依據 profiler 得出的機制歸因,不是這份官方文件直接給出的效能保證。
cuda-oxide 本身也仍屬早期實驗性編譯器。其官方專案明載目前為 alpha、應預期錯誤、不完整功能與 API 破壞性變更;同時確認它是把 Rust 核心編譯至 PTX 的自訂 rustc 後端:https://github.com/NVlabs/cuda-oxide。論文期間發現 scoped atomic load/store 無法在實際核心建置模式使用,修補已合併上游。這應視為已修復的實作缺陷,而不是 Rust 語言本身的永久限制。
容易忽略的限制
數字只來自 sm_120 Blackwell 架構的 RTX 5070 Ti 與 RTX 5060,不能直接外推至資料中心 GPU、舊世代硬體或其他廠牌。所有實作亦由同一作者完成,而作者自陳 CUDA 經驗多於 Triton;固定相同資料布局雖讓比較公平,卻也阻止 Triton改用更適合自身模型的演算法與結構。
此外,九組所謂「真實深度」資料實際來自 TartanAir 模擬環境,具有精確 ground truth,沒有實體感測器的雜訊、掉點與校正誤差。Rust 修正後仍有中位數約 1.21 倍的 allocate 殘餘差距,作者明確表示尚未歸因。這些都是來源事實;把結論擴大成「Rust 已取代 CUDA」或「Triton 不適合非 AI 工作」則沒有資料支持。
對台灣 GPU、機器人與邊緣視覺團隊的意義
作者觀點:工具鏈選擇應以核心形狀分區,而非要求整個產品押注單一語言。矩陣、卷積與規則累加可利用 Triton 的開發效率;雜湊插入、稀疏配置、動態佇列與高度爭用原子操作,則應保留 CUDA C++,或在充分剖析後評估 Rust。混合工具鏈增加建置與維運成本,卻可能比勉強讓一種抽象涵蓋所有核心更實際。
對自主移動機器人、機械手臂、數位孿生及工業 3D 視覺而言,正確性閘門也應與效能測試綁在一起:比對配置區塊數、回報探測上限與池耗盡、以已知幾何檢查表面誤差,並把配置、更新與資料搬移分段計時。論文的程式、產表腳本與原始 CSV 已公開,可用來檢查實驗如何生成數字:https://github.com/realitymatrix/what-irregularity-costs
接下來應觀察什麼
- Triton 是否為 atomic_cas 增加遮罩,或提供可表達逐 lane 提前退出的替代模型。
- 相同測試能否在 Hopper、Ada 或資料中心 GPU 重現排序,以及絕對倍率是否明顯改變。
- cuda-oxide 的原子操作修補、CI 覆蓋與 API 穩定度,能否從研究可用走向正式產品可維護。
- 在實體感測器與更高雜湊負載下,Triton 的資料遺失門檻及 Rust 未解殘差如何變化。
這篇研究最值得帶走的結論,不是 CUDA C++ 永遠最快,而是「抽象的邊界就是成本的起點」。當基準離開整齊矩陣,真正需要量測的已不只是 FLOPS,而是工具是否能精確表達提前退出、同步範圍、爭用與錯誤回報。