連英偉達都開始抄作業，國內 AI 的野路子有點猛

咱們都知道，英偉達在硬件上賣 GPU 算力，軟件方面手握 CUDA 生態，都快把顯卡和 GPU 計算這門 “給 AI 賣鏟子” 的生意給壟斷了。

可以說在 GPU 編程這件事上，CUDA 幾乎就是 “版本答案”，以至于之前行業里形成了一種共識：想要真正發揮顯卡的性能，基本繞不開 CUDA，自然也就繞不開老黃的顯卡。

但是前一陣老黃發了個 GPU 編程的新技術，居然是反過來向中國 “取經”？

事情是這樣的，上個月初英偉達毫無征兆地推出了一個叫 CUDA Tile 的 GPU 編程語言，直接讓圈里人炸鍋了。因為官方是這么形容這次更新的，“這是自 2006 年 CUDA 發布以來，最大的一次進步”。

那為什么說老黃這個技術是在向中國取經呢？因為 2024 年的時候，剛有一個中國團隊開發過和 CUDA Tile 思路非常相似的 GPU 編程語言 “TileLang”。

并且前段時間 DeepSeek 發布 V3.2 的時候，真的同步發布了兩個版本 —— 一個 CUDA 語言構建的版本，一個 TileLang 語言構建的版本。

曾經 DeepSeek 只發布 CUDA 語言構建的版本，甚至為了榨干硬件算力，不少代碼用的還是英偉達專有的 PTX 匯編語言，導致過往的 DeepSeek 模型與老黃的運算卡深度綁定。

但是中國團隊開發的 TileLang 編程語言不挑運算卡，所以 DeepSeek V3.2 可以很輕松的部署在華為昇騰等國產運算平臺上。

所以事情開始變得有意思起來了。。。由于 TileLang 發布、DeepSeek 使用 TileLang、CUDA Tile 發布這幾個事兒前后腳挨著，托尼就在網上看到有種說法：老黃是感受到國內廠商的威脅了，這才開始 “抄作業”。

那么真相到底是不是這樣呢？一個 GPU 編程語言，就能動搖英偉達的根基了？

真不怨網友們唱衰老黃，托尼身邊熟悉 GPU 編程的人，最近幾年也都在吐槽一件事：現在的 CUDA，好像有點跟不上 AI 時代的節奏了。

過去 CUDA 之所以讓大家覺得效率爆炸，靠的其實是一套非常省事的工作方式，叫 SIMT（單指令多線程）。簡單說，就是 GPU 發一條指令，就能讓一大堆線程同時行動。

就拿 FPS 游戲舉例，當對面朝你扔了一顆閃光彈，GPU 干的活很簡單：對著屏幕上幾百萬個像素，重復同一套操作 —— 算亮度、填顏色，讓畫面瞬間變得“晃眼”。

渲染管線（頂點→像素）著色過程

其實不光是“像素著色”這種活兒，實際上在圖形渲染和通用計算的時代，交給 GPU 來處理的工作，都可以理解為類似的流水線作業：任務單一、步驟固定、數量還多。GPU 最擅長的正是這種一件規則完全一樣的事，同時做成千上萬遍。

但到了 AI 時代，GPU 發現天塌了，因為 AI 推理不像渲染那樣，一眼就知道接下來要干啥，而是要“摸著石頭過河” —— 后面的計算，往往得等前一步的結果出來，才能決定。

就拿大家常用的 OCR、語音識別、AI 問答，或者生成式模型來說：有的字一眼就認出來了，有的得反復確認；有的語音一句話就結束，有的要來回糾錯；有時候你只是中間換了個提示詞，模型后面的計算流程就徹底變了。

問題在于，SIMT 這套機制一點沒變，GPU 還是那套老習慣：一組線程，必須聽同一條指令，按同一個節奏往前走。所以就出現了線程之間彼此看眼色，快的等慢的、算完的等沒算完的，GPU 的效率就被浪費了。

而且還有一個更現實的問題：在 AI 推理中，很多運算的中間結果會被反復使用，但 CUDA 的執行模型并不關心數據復用，一輪線程執行結束，結果就會被寫回全局內存。

所以 CUDA 編程被吐槽難寫、跟不上時代，就在于寫代碼的人不光要懂算法，還要把線程的 “組織架構和分工” 排明白：哪些線程負責搬數據，哪些線程負責計算，什么時候同步，全都要程序員手搓完成。。。搞不好就造成 GPU 的效率打折扣。

其實英偉達很早就意識到了 CUDA 對 AI 不友好。所以在 2014 年的時候，老黃又端出來了一個叫 “cuDNN” 的玩意，它的思路非常直接 ——

既然很多 AI 算子大家用 CUDA 從頭寫，既復雜又容易出錯，所以干脆先讓英偉達的工程師們，把最常用的幾種計算，比如卷積、矩陣乘，全部提前寫好，再封裝成開發者可以直接調用的內核。

換句話說，cuDNN 就是一些 AI 推理模型的 “工業預制菜”，程序員只管點菜，GPU 內部怎么跑、怎么調，交給 cuDNN 自己處理。

所以 cuDNN 的缺點也很明顯：AI 研究員永遠只能點菜單上的菜。一旦模型里出現了新算子，對于 cuDNN 來說考題就超綱了。。。這時候 AI 研究員還是要回到 CUDA，重新和線程、內存、同步這些底層細節打交道。

而 TileLang 的出現，改變的正是這一點。它并不是再多做一份“預制菜”，而是直接把原來程序員要干的那一大堆“調度”雜活 —— 線程怎么分、數據怎么復用、什么時候同步，全都接了過去。

左圖-從整體看任務拆分，右圖-從局部看線程如何逐元素執行；CUDA 的線程思維 thread-level 屬于右邊這種～

這下開發者只需要告訴 TileLang：想做什么計算，用哪些數據，怎么算。至于這些計算怎么映射到具體的 GPU 硬件和線程上，統統交給編譯器去完成。

結果就是，即使是 cuDNN 沒覆蓋的新算子，程序員也不必再回到手寫 CUDA 的老路上。代碼更少，邏輯更清楚，性能也更容易跑出來。

根據官方示例，在一些算子開發中，TileLang 可以把 CUDA/C++ 的代碼量從 500 多行壓縮到約 80 行，同時性能反而提升了約 30%。過去“又累又慢”的 GPU 編程，這一次在開發效率和運行性能上，都吃到了 TileLang 帶來的紅利。

在 H100 顯卡上，計算 MLA 算子，TileLang 性能接近目前最快的 FlashMLA

更關鍵的是，TileLang 這類工具還順手驗證了一件以前大家不太確定的事 ——

只要編程語言能把計算邏輯表達清楚、編譯器足夠聰明，不用手寫 CUDA 這種底層代碼，GPU 同樣可以跑得很快。

事情發展到這里，英偉達其實已經很難繼續保持觀望了。

所以去年 12 月英偉達親自下場，推出了 CUDA Tile，你可以理解為英偉達官方版本的 TileLang。

這個動作傳遞的信號其實非常明確：與其讓開發者依賴第三方 GPU 語言，慢慢摸索 “GPU 還能怎么用”，不如這件事，英偉達自己來干。

作為英偉達的第一方 GPU 語言，CUDA Tile 掌握著最短最直接的優化路徑。在可預見的未來，如果想在英偉達 GPU 上，榨干每一分性能，那 CUDA Tile 大概率會是最省事、也最穩妥的選擇。

要說 CUDA Tile 一出來，TileLang 就沒戲了？那也不至于。

因為 CUDA Tile 再強，也有一個前提：你打算一直用老黃的顯卡。

而 TileLang 的最大價值，就在于它不被任何一家硬件廠商綁定 ——

因為過去大家買顯卡，關注的是：“這張卡的 CUDA 生態成不成熟，寫編程時的現成工具多不多？”

但當 TileLang 這種寫法越來越常用之后，問題可能會變成：“我現在這套用 TielLang 寫的代碼，換一張卡還能不能繼續跑？”

無論是 AMD 的 GPU、谷歌的 TPU，還是國產 AI 芯片，只要 TileLang 跑得通，那么 AI 模型和算子，就不再必須按照 CUDA 的方式來寫，開發者就不用為了換平臺，把模型和算子推倒重寫一遍，相當于所有訓練模型的人都擁有了 “自己挑鏟子” 的能力。

類似的故事我們并不是第一次見到。

游戲市場里， DirectX 12 已經和 Windows 深度綁定，可以在 Windows 平臺上性能發揮到極致、工具鏈和優化經驗也非常成熟，但這并沒有阻止 Vulkan 這樣更開放、更跨平臺的技術，逐步分走開發者的聲量和生態位置。

開發者用腳投票的結果已經證明了，性能并不是唯一標準，開發者為了不被單一廠商的技術路線“卡脖子”，有時候也會主動選擇更開放的技術路線。。。

撰文：Levi

編輯：米羅 & 粿條 &面線

美編：煥妍

圖片、資料來源：

英偉達官網

mashdigi.com

tilelang.com

medium.com

csdn @Wanderer001 —— CTC（Connectionist Temporal Classification）介紹