晶片層 開放閱讀

SIMT

Single Instruction, Multiple Threads

概念 ID
single-instruction-multiple-threads
更新時間
2026-05-29
來源數量
待補

SIMT 架構

摘要

單指令多執行緒(Single Instruction, Multiple Threads,SIMT)是當代高效能運算與人工智慧訓練的核心並行架構模型。它並非一個孤立的概念,而是融合了早期向量處理思想、SIMD資料並行模式與現代多執行緒硬體管理技術的集大成者。本研究報告旨在提供一份超過萬字的深度技術產業分析,全面解構SIMT從底層執行哲學到上層生態系統的每一個關鍵層面。我們將摒棄抽象的技術讚美,深入其執行模型的微觀運作、記憶體層次的設計權衡、排程策略的效率博弈、分支發散的硬體與軟體協同處理、以及其如何建置起人工智慧計算的堅固基座。本報告將沿著一條邏輯嚴密的主線,從歷史源流出發,歷經執行模型、記憶體系統、控制流挑戰、硬體微架構實現、程式設計生態建置、到人工智慧負載的適應性分析、競合對比、製造工藝影響以及未來演進方向等十五個核心維度,為您建立一個關於GPU計算基座的、多維度的系統性認知。

1. 引言:計算範式的轉折與SIMT架構的必然性

過去二十年,全球計算版圖經歷了一場足以載入史冊的範式大遷移——從以CPU為中心的順序標量處理,邁向以GPU/GPGPU為中心的大規模資料並行處理。這場變革的深層動力並非單純來自半導體工藝的線性演進,而是源於計算需求側的結構性突變:資料級並行(DLP)需求的爆炸式增長,尤其是在2012年AlexNet橫空出世後,以深度學習為代表的人工智慧工作負載,對浮點運算吞吐量提出了無止境的渴求。

傳統的單核乃至多核CPU架構,其設計哲學根植於最佳化單一執行緒的執行延遲。龐大的快取層級、複雜的亂序執行邏輯、高精度的分支預測器,這一切都是為了讓一個序列程式流盡可能快地走完。然而,當面對一個需要為千億引數模型執行數萬億次矩陣乘加運算的任務時,這種“延遲導向”的策略在經濟上迅速失效。CPU核心數量的有限擴充套件性與其每個核心高昂的電晶體開銷,在面對“易並行”的稠密線性代數計算時,猶如用一支精雕細琢的鋼筆去繪製一幅巨型壁畫,雖有精度卻無效率。

正是在這種歷史張力之下,一種名為“單指令多執行緒”(SIMT)的架構哲學從嵌入式圖形處理的特定領域走出,並迅速演化為通用平行計算的支配性範式。SIMT並非橫空出世的天才發明,它有著清晰的技術譜系。其思想源頭可追溯至1970年代的向量處理器(如Cray-1),向量機通過深度流水化的功能單元對連續資料向量進行操作;更直接的先驅是單指令多資料(SIMD)指令集擴充套件,如Intel的MMX/SSE/AVX,它們允許一條指令同時處理暫存器中打包的多個短資料。然而,SIMT實現了一次決定性的飛躍:它以完全標量的程式設計介面,隱藏了向量的寬度和執行掩碼,並通過硬體管理的海量輕量級執行緒,實現了對資料級並行和執行緒級並行的雙重榨取。

2006年,NVIDIA推出G80圖形核心及其配套的CUDA程式設計模型,是SIMT走向產業化的“奠基時刻”。G80並非只是另一款更快的GPU,它通過統一著色器架構將圖形管線的頂點、畫素、幾何處理單元合併為一個通用的流多處理器(Streaming Multiprocessor,SM)陣列,並首次向開發者暴露了一個類C的程式設計環境。其底層硬體哲學便是SIMT:程式設計師為單個執行緒編寫標量程式碼,硬體在執行時將此程式碼例項化到同一計算核心上的數百個執行緒,並以32個執行緒一組的“執行緒束”(Warp)為單位,以鎖步方式執行。這一轉變使得並行程式設計的抽象成本大幅降低,將科學家和工程師從手動畫分向量長度、管理資料對齊和編排顯式資料並行的繁瑣勞動中解放出來。

今天,SIMT架構的影響已遠超GPU本身。無論是訓練一個擁有萬億級引數的GPT-4多模態模型,還是驅動自動駕駛汽車在毫秒級延遲內完成多感測器融合感知,抑或是在邊緣裝置上進行即時的Stable Diffusion推論,其算力的根基無一例外地建立在數百萬個遵循SIMT模型的CUDA Core、Tensor Core或Stream Processor的密集陣列之上。從Google的TPU到新興的專用AI晶片,其脈動陣列或向量處理單元中依然能窺見SIMT執行模型的影子。

本研究報告試圖穿透表面的規格引數與營銷話術,為讀者揭示SIMT架構的真正技術精髓:它如何用一條統一的指令流水線驅動成千上萬條獨立的資料流?其硬體排程器如何在納秒級時間內完成執行緒的零開銷上下文切換以掩藏記憶體訪問延遲?面對程式中普遍存在的if-else分支,它如何通過棧與執行掩碼的協同巧妙化解“分支發散”的效能災難?以及,最關鍵的是,為什麼這種看似為圖形渲染設計的模型,最終能夠在人工智慧時代構築起一個市值數萬億美元的、近乎堅不可摧的軟體生態壁壘?我們需要從最直觀的定義開始,逐步深入其微觀世界。

2. 三秒穿透本質:SIMT架構定義的精髓解構

若要在紛繁的技術術語中提煉SIMT的靈魂,其核心定義可表述為:一種平行計算架構模型,其中一個指令單元在給定的時鐘週期內,將同一條指令廣播給由多個獨立執行緒組成的執行組,組內的每個執行緒持有各自的程式計數器、暫存器檔案和本地資料,在同一時刻對不同的資料元素執行完全相同的操作。 這種執行模式的最直觀類比,是一位軍事教官向一整排士兵下達“舉槍”的統一口令。每一位士兵接收到的是同一個指令,但每位士兵舉起的是自己手中的、可能是不同型號、不同序列號的那支槍。動作是同步的,但操作的物件是個性化的。

這一模型與經典SIMD(如SSE/AVX)的根本區別在於“執行緒”的抽象層級。SIMD指令直接操作固定寬度的向量暫存器。例如,一條 _mm256_add_ps 指令會同時將兩個256位暫存器中的各8個單精度浮點數進行相加。這要求程式設計師顯式地知道向量寬度,並精心組織資料以滿足對齊和分塊要求。一旦向量寬度改變(例如從AVX2 256位升級到AVX-512 512位),舊的程式碼需要大量重寫以充分利用新硬體。

SIMT通過引入“執行緒”這一程式設計抽象,徹底透明化了硬體向量寬度。程式設計師仍以標量方式為單個執行緒編寫核心函式(kernel)。例如,執行 c = a + b,假定這只是一個執行緒對自己的兩個私有變數進行相加。當程式設計師請求啟動數千個這樣的執行緒時,硬體負責將這些執行緒分組打包。在NVIDIA硬體中,32個執行緒被組合成一個執行緒束,並在一個SM內的CUDA核心上以鎖步方式執行。從程式設計師視角看,他是為每一個數據點單獨編寫了一個執行緒,硬體向量寬度對他不可見。硬體保證這個標量加法的操作,會以32個執行緒一組的方式被廣播到執行單元陣列。這種標量程式設計,向量執行的哲學,是SIMT降低並行程式設計門檻、獲得巨大軟體生態成功的秘訣。

更深一層,SIMT架構為每個執行緒維護了獨立的執行上下文,這在傳統SIMD中是不存在的。每個SIMT執行緒擁有自己的通用暫存器堆(Register File)空間,用於存放其私有的區域性變數和中間結果。指令單元以執行緒束為粒度獲取指令,但其源運算元和目標運算元的暫存器地址,對於束內每個執行緒是不同的。這意味著,硬體必須維護一個巨大的、多埠的暫存器檔案,並能同時為執行緒束內的所有執行緒提供讀/寫服務。這構成了SIMT處理器功耗和麵積的主要部分,也是其區別於CPU的顯著特徵——CPU核心維護少量顯式執行緒(超執行緒)的龐大狀態,而GPU的SM則維護大量細粒度硬體執行緒的緊湊狀態。

從產業視角看,該模型由NVIDIA在其CUDA使能的GPU中首次系統化實現並推向市場極致。它對開發者和市場的核心承諾,構成了其商業模式大廈的基石:你像為單個數據點編寫序列程式一樣編寫GPU核心,編譯器和硬體負責將你的邏輯無縫對映到晶片上數千個並行物理核心上。 這一看似簡單的承諾,讓原本需要高深平行計算領域知識的GPGPU開發,向數百萬懂C/C++的普通程式設計師敞開了大門,從而催生了繁榮的CUDA生態,並最終使GPU成為人工智慧時代的通用計算引擎。理解這一點,是理解輝達為何能從一個圖形晶片商蛻變為AI算力霸主的邏輯起點。

3. 產業邏輯:為何SIMT定義了人工智慧算力基座

從產業演進的宏觀視角審視,SIMT架構在人工智慧領域的統治地位,是一場技術稟賦、經濟激勵和生態網路效應共同作用的完美風暴。它並非簡單地“因為快所以被選中”,而是其內在的執行模型、記憶體體系與人工智慧核心計算模式的數學特徵之間,存在著一種近乎“天作之合”的深層同構性。

人工智慧大型模型(無論是GPT系列的大型語言模型、Llama系列的開源模型,還是Sora等擴散生成模型)的訓練和推論,其計算核心可以抽象為極度稠密、高維度且“易並行”的張量運算。以深度學習中最核心的運算元——通用矩陣乘法(GEMM)為例,計算 C = A \times B。在這個運算中,輸出矩陣 C 中的每一個元素 C_{i,j} 都是矩陣 A 的第 i 行與矩陣 B 的第 j 列的點積。計算這一萬個甚至一百萬個 C_{i,j} 的單元之間,不存在任何資料依賴關係,它們可以完美地同時進行計算。這種計算形態,被稱為“令人尷尬的並行”。

SIMT架構的天才之處在於,它為這種“易並行”模式提供了一種極其廉價且高效的實現方案。它將一個由數百萬個獨立計算組成的全域性問題,分解為層次化的執行緒網格:

  1. 執行緒塊網格:整個計算任務被劃分給一個一維、二維或三維的執行緒塊(Thread Block)網格。每個執行緒塊被排程到任意一個流多處理器(SM)上獨立執行。執行緒塊之間完全解耦,可以以任意順序在任何空閒的SM上執行,這賦予了架構極強的可擴充套件性——從1個SM到100個SM,程式碼不變,效能線性擴充套件。
  2. 執行緒束執行:在一個執行緒塊內部,執行緒被進一步分組為執行緒束(NVIDIA是32執行緒)。執行緒束是SM內部指令發射的基本單元。當執行GEMM時,SM的指令排程器將會把一條“融合乘加”指令廣播給執行緒束內的所有32個執行緒。這32個執行緒可能各自從A矩陣的不同行和B矩陣的不同列的同一個索引位置獲取資料,執行乘法後再累加到各自的區域性和中。
  3. 單執行緒向量化:每個SIMT執行緒內部,還可以通過載入128位(4個浮點數)的float4向量進一步榨取訪存頻寬,實現類似SIMD的效果但又不改變執行緒級的標量程式設計模型。

這種“一個執行緒負責一個或幾個結果元素,所有執行緒同步推進”的執行模式,對於GEMM類運算,意味著極高的理論峰值得以實現。硬體只需設計足夠寬的指令發射寬度(例如同時發射給4個執行緒束)和足夠多的執行單元,就能將計算吞吐量線性堆疊上去。NVIDIA引入的Tensor Core更是將這一模式推向量子化:一個執行緒束內的32個執行緒,協同向一個專用硬體單元提供一個4x4x4矩陣乘法的分片(Tile),Tensor Core在單週期內完成這個小型矩陣乘法,實現了遠超標量CUDA Core的吞吐量。從產業看,這意味著,誰能推出擁有更多SM、更寬Tensor Core並有效利用SIMT排程這一切的晶片,誰就能定義大型模型訓練的“時間刻度”,因為訓練時間與可用算力成反比。

而對於軟體產業生態而言,SIMT創造的價值是對複雜性的完美封裝。PyTorch等頂層架構的開發者,不需要關心底層的執行緒束大小是32還是64。當他們呼叫 torch.nn.Linear 時,這一操作被層層下放至cuBLAS等庫,再由庫生成最優的SIMT執行緒網格配置。這個分層解耦的軟體棧,使得演算法創新者可以專注於模型架構的數學設計,而無需觸碰複雜的平行計算原語。這種分工帶來了效率的極大釋放:成千上萬的人工智慧演算法工程師和科學家,雖然不懂SIMT為何物,卻每天都在通過PyTorch間接呼叫它,創造了龐大的需求。

這種供需兩側的互鎖效應,是SIMT生態壁壘的根本來源。NVIDIA憑藉CUDA將SIMT硬體的價值無縫傳遞到每一個Python指令碼和每一個深度學習運算元中,形成了近二十年的積累。任何一個新入局的AI晶片挑戰者,即便其硬體在紙面上的理論算力更高,若無法在其上建置出同等豐富、同等最佳化程度的,無縫對接PyTorch生態的SIMT或類SIMT軟體棧,便難以撼動這個由數百萬開發者和海量程式碼庫構築的、沉沒成本極高的產業格局。因此,可以說,SIMT架構不僅承載了人工智慧的算力,更通過其程式設計模型定義了當今人工智慧研發的工程範式。

4. 歷史演化:從向量機、SIMD到現代SIMT的螺旋上升

理解SIMT,必須將其置於平行計算架構數百年的宏觀演進中考察。這條主線貫穿了從早期向量處理機到現代通用GPU的螺旋上升史,是一個不斷在掩藏延遲、提升峰值吞吐量和改善可程式設計性之間尋求更優解的過程。

第一階段:向量處理機的黎明(1970s) 1976年,Cray-1超級計算機的誕生標誌著向量處理架構的成熟。Cray-1的核心是其向量處理器,它使用深度流水線(例如12級浮點乘法流水線)對儲存在向量暫存器中的密集資料陣列進行操作。一條“向量加”指令可以啟動流水線,並在後續每個時鐘週期流出一個結果,經過一定啟動延遲後,實現每個週期一個結果的吞吐量。這是資料並行最原始也最強大的形式之一。然而,向量機對編譯器高度依賴,向量化是程式設計師/編譯器必須顯式完成的工作,對含有條件判斷的稀疏操作處理笨拙,且其高昂的專有硬體成本限制了它的普及。

第二階段:SIMD指令集擴充套件的普及(1990s末-至今) 隨著多媒體和科學計算在個人電腦上的興起,為通用標量處理器增加SIMD功能單元成為經濟之舉。Intel的MMX(1997年)、SSE(1999年)系列及後續的AVX、AVX-512是這一路線的代表。SIMD指令讓一個窄位寬(如64位)的CPU核心能在一個寬向量暫存器(如128/256/512位)上打包操作多個短整數或小數。這是“單指令多資料”在消費級硬體上的巨大成功。但是,其侷限也很顯著:向量寬度固化在指令集(ISA)中,軟體升級換代困難;程式設計師必須進行復雜的內聯彙編或Intrinsics程式設計,可移植性差;當資料不能完美對齊或資料間存在條件分支時,效能會急劇下降。

第三階段:SIMT的誕生與GPGPU融合(2000s初-中) 圖形處理的特殊需求催化了SIMT的雛形。在畫素著色階段,一個三角形光柵化後會生成成千上萬個獨立的畫素碎片,對它們進行的紋理取樣和光照計算是高度並行的,且每個畫素需要維護自己的顏色和深度等狀態。這天然地催生了一種在每個“畫素執行緒”上執行相同著色器程式,但資料各自獨立的架構,這便是GPU的根源。早期GPU的頂點和畫素管線是分離的,各有其專用硬體。

2006年NVIDIA G80架構是一個分水嶺,它實現了“統一著色器”模型。硬體上將頂點、畫素、幾何著色器統一為執行通用程式的流多處理器(SM)。邏輯上,NVIDIA將這種執行模型正式抽象為SIMT。G80的每一個SM可管理多達768個活動執行緒,並以24個執行緒的執行緒束(在後續架構中演變為32)為單位進行排程。這帶來兩個革命性變化:

  1. 細粒度多執行緒:當一個執行緒束因紋理/記憶體訪問而停滯時,硬體排程器可在零開銷下切換到另一個就緒的執行緒束,用計算隱藏訪存延遲。
  2. 標量ISA:儘管以執行緒束鎖步方式執行,但其指令集是標量化的。開發者用“單資料”模式思考,硬體完成“多資料”的廣播。這為CUDA C的出現鋪平了道路。

第四階段:AI算力專用化與持續精進(2010s末-至今) SIMT架構的靈活性使其能不斷吸收新的專用計算單元。2017年,NVIDIA Volta架構引入了第一代Tensor Core,專門用於加速矩陣乘加運算。這並非對SIMT模型的背離,而是一種增強。執行緒束的32個執行緒現在可以協同工作,共同向一個Tensor Core提交一個4x4x4的矩陣分片任務。Tensor Core本身是高度並行化的脈動陣列或類似結構,但其任務的分發、編排和結果寫回依然在SIMT排程模型的統一架構下。這展現了SIMT作為一個“元架構”的強大相容性。與此同時,AMD的GCN及後續RDNA架構也發展出其稱為“波前”(Wavefront)的SIMT變體,執行緒組大小通常為64。兩者殊途同歸,共同印證了SIMT作為大規模資料平行計算最優解的普適性。

5. 微觀執行模型:執行緒束鎖步、排程與掩碼機制

要理解SIMT的效能之源與瓶頸之根,必須深入其微觀執行模型。這無關CUDA的程式設計抽象,而是關乎硬體在每一個時鐘週期內,如何在指令流水線、暫存器堆和執行單元之間精確地搬運和運算元據。其運作精髓體現在三個核心機制上:執行緒束鎖步、細粒度多執行緒排程和執行掩碼。

執行緒束鎖步執行:並行性的物理根源 SIMT架構的並行性並非靠亂序超標量邏輯實現,而是靠“單指令廣播”的鎖步執行。在一個流多處理器內部,指令獲取和發射單元以執行緒束為粒度工作。當一個執行緒束經排程器選中並執行“載入-儲存”或“乘加”指令時,該指令的操作碼被同時廣播到一組執行單元。每個執行單元都連線到一個特定的執行緒,並從該執行緒的私有暫存器檔案中讀取源運算元。關鍵在於,這32個執行緒(以NVIDIA為例)共用一個程式計數器,但擁有獨立的暫存器空間。這要求SM內部的暫存器檔案必須是高度分塊、多埠設計的,能夠同時支援32路獨立的讀寫操作。這種設計放棄了單執行緒的高主頻,換取了極致的吞吐面積比和能效比。

細粒度多執行緒排程:隱藏延遲的唯一法門 GPU與CPU一樣,都要面對記憶體訪問的長延遲(通常數百個時鐘週期)。CPU的策略是建置龐大的私有快取(如L2、L3)和使用複雜的亂序執行,試圖減少平均延遲。這是一種“延遲導向”的設計。SIMT採取的策略則完全不同,是一種“吞吐導向”的設計:當一個執行緒束(Warp A)發出一個全域性記憶體讀取指令後,它將在數十到數百個週期內無法繼續執行。此時,硬體執行緒束排程器絕不空轉等待,而是立即將Warp A移出發射槽,並從同一SM上駐留的其餘就緒執行緒束中選出一個(如Warp B)發射其下一條指令。只要SM上始終有足夠多的活動執行緒束可以執行非訪存指令(如算術運算),訪存延遲就會被完美地隱藏在持續的計算流之下。這種纖程級的零開銷上下文切換,是SIMT處理器能低成本利用高頻寬DRAM(如HBM)的秘訣。如果一個SM上的活動執行緒束數量不足,計算管線就會因飢餓而氣蝕,效能急劇下降。因此,“佔有率”,即SM上活動執行緒束數量與理論最大支援數量的比值,是最佳化CUDA程式效能的首要指標。

執行掩碼與分支發散:統一指令流的分化處理 當程式中遇到條件語句(如 if(threadIdx.x % 2) { path_a; } else { path_b; } )時,鎖步執行的悖論便出現了。若執行緒束內部分執行緒走 path_a,另一部分走 path_b,由於它們共享一個程式計數器,無法同時執行兩條路徑。SIMT的硬體解決方案是“執行掩碼”。當遇到分支點,排程器會設定一個掩碼,首先將執行 path_a 的執行緒位設為1,而將要走 path_b 的執行緒位暫時遮蔽為0。然後該執行緒束全體單步遍歷 path_a,被遮蔽的執行緒處於休眠狀態,但其指令仍流過管線(或被丟棄)。之後,掩碼翻轉,再執行 path_b。最終,兩條路徑都執行完畢,掩碼重設為全1,執行緒束再次匯合,繼續同步向下執行。這個過程稱為“分支發散”,代價是損失了執行 path_apath_b 的並行度,執行時間等於兩者之和。現代架構還利用“獨立執行緒排程”等技術,允許束內執行緒擁有自己的程式計數器,實現更細粒度的併發,但最終匯合機制仍是管好分支發散效能代價的核心。

6. 記憶體層次與矩陣:頻寬、延遲與資料區域性性的博弈

如果說SIMT的執行模型是引擎,那麼其精心設計的多級記憶體層次就是為這個引擎輸送燃料的高效油路系統。此係統設計的核心訴求只有一個:在每一筆運算消耗一個單位能量、佔用一個單元面積的情況下,從記憶體中供給儘可能多的資料。這是一個關於頻寬、延遲和訪存區域性性的永恆博弈。

全域性記憶體與L2快取:大容量、高頻寬的共享底座 這是GPU板載的HBM/GDDR記憶體。它是離計算核心最遠、容量最大、頻寬最高但延遲也最高的儲存層級。A100 GPU的HBM2e頻寬可達2TB/s,H100的HBM3頻寬高達3TB/s以上。如此驚人的頻寬並非為某個執行緒獨享,而是為了滿足一個圖形核心上所有SM、所有活動執行緒束不斷髮出的合併訪問請求。SIMT程式設計模型中一個核心最佳化概念是“合併訪問”:當同一個執行緒束內的32個執行緒所訪問的全域性記憶體地址,落入同一個32位元組、128位元組對齊的記憶體段(segment)內時,這些分散的請求會被記憶體控制器合併為一個寬匯流排事務。一次事務取回一個連續的資料塊,然後自動分發給各個執行緒。如果一個執行緒束的訪問模式是雜湊的、非合併的,那麼記憶體子系統將不得不發起多個獨立事務,有效頻寬將斷崖式下降。因此,資料的組織和執行緒與資料的對映方式,是獲得高頻寬的決定性因素。

共享記憶體與L1快取:程式設計師手控的低延遲“便籤” 這是GPU片上最寶貴的儲存資源。它是一個由程式設計師顯式管理的、極低延遲的便籤式儲存器(Scratchpad RAM),物理上緊鄰SM內的ALU。其重要性再怎麼強調也不為過。在一個矩陣乘法分片(Tiling)策略中,程式設計師會將A和B矩陣的小分片,通過執行緒束間的協同載入,從全域性記憶體搬運至共享記憶體中。之後,整個執行緒塊內的所有執行緒便可以極低的延遲和極高的頻寬,反覆讀取共享記憶體中的分片資料進行乘加運算,而無需再次觸碰高延遲的全域性記憶體。這種“計算前,先協同搬運資料至近端;在近端資料上反覆迭代計算”的模式,是充分釋放GPU算力的標準範式。共享記憶體完全由軟體控制,沒有快取的不確定性,是高手調優的舞臺。

暫存器檔案:執行緒私有計算速度的極限 這是GPU記憶體層次中速度最快、位於金字塔頂端的儲存。每個SIMT執行緒的所有標量區域性變數和中間計算結果,都存放於其私有的暫存器中。GPU對於暫存器消耗極度敏感。一個SM上總的暫存器檔案大小是固定的(如65536個32位暫存器)。如果每個執行緒佔用過多暫存器,SM能同時容納的活動執行緒束數量(即佔有率)就會降低,從而削弱其用切換隱藏延遲的能力,這被稱為“暫存器壓力”。編譯器在編譯CUDA核心時,會進行復雜的暫存器分配最佳化,試圖在減少單執行緒暫存器佔用(提升並行度)和避免將變數溢位(spill)到緩慢的本地記憶體之間找到最佳平衡點。本地記憶體其實是全域性記憶體或L1/L2快取中的一片私有區域,一旦發生溢位,效能將受到嚴重打擊。

整個記憶體體系構成了一個嚴密的金字塔:從TB/s級、高延遲的HBM,到由軟體協同搬運、低延遲的共享記憶體,再到每個執行緒零延遲可用的暫存器檔案。SIMT架構的效能天花板,最終往往不取決於理論FLOPS,而取決於資料在上述各級儲存層級間移動的演算法效率。

7. 控制流與分支發散:從鎖步困境到現代最佳化策略

控制流,尤其是條件分支,是SIMT架構需克服的根本性挑戰。它源於鎖步執行模型與程式邏輯必然存在的多樣性之間的固有張力。當一個執行緒束內不同的執行緒執行不同的程式碼路徑時,便產生了“分支發散”,導致並行效率折損。對這一問題從被動容忍到主動最佳化的演進,構成了GPU架構設計史的另一條主線。

發散的本質與代價模型 分支發散問題在SIMT模型中表現為執行緒束內執行緒計算的序列化。考慮一個簡單的if-else語句,假設執行緒束中有x個執行緒進入if分支,y個執行緒進入else分支(x+y=32)。那麼,這個if-else結構的執行時間將與x+y成正比,而不是為Max(x, y)。在最壞情況下,如果分支是完全隨機的,一個執行緒束內每個執行緒都走獨一無二的路徑,理論上效能會降至完全無並行的序列化。這種發散效能損失不僅是此分支本身的損失,更嚴重的是,這32個執行緒在發散後佔用了該執行緒束完整的執行時隙,但仍然要消耗暫存器和其他資源,阻塞了它們在各自分支路徑上並行執行的可能。

架構演進:從棧式掩碼到獨立執行緒排程 早期SIMT架構使用一個簡單的實現方式:硬體內部維護一個細粒度的執行掩碼棧。每進入一個巢狀的分支,當前活動掩碼被壓入棧中,並生成新的掩碼。退出分支時,棧被彈出並恢復掩碼,確保執行緒重新匯合。這種方法在硬體控制上是直接的,但程式設計師對發散的控制力不強。

NVIDIA在Volta架構及以後引入了“獨立執行緒排程”模式。這改變了鎖步的剛性限制,允許執行緒束內所有執行緒擁有獨立的程式計數器和呼叫棧。硬體排程器現在可以以單個執行緒的粒度交錯執行不同路徑上的指令。一個執行緒在執行if塊的程式碼時,另一個已經可以開始執行else塊。這極大地提高了某些動態發散的複雜負載(如光線追蹤、某些圖演算法)的利用率。然而,這並不意味著分支發散的代價為零,因為它增加了排程複雜性和對片上指令快取、排程的壓力,且最終執行緒在需要通訊或匯合時仍需同步等待。對於主流AI負載,由於稠密計算的分支極為規整(常以執行緒塊或完整執行緒束為粒度分界),傳統鎖步發散損失很小,獨立執行緒排程的價值更多地體現在通用性上。

演算法與編譯器協同最佳化 對抗分支發散,最重要的武器在軟體層面。最佳化策略包括:

  1. 資料重排:在預處理階段重新組織資料,使得同一執行緒束內的執行緒傾向於處理在邏輯上會走相同分支的資料。例如,在粒子分類或光線追蹤中,按材質或包圍盒對資料進行排序。
  2. 分支規整化:編譯器或程式設計師將原來可能隨機發散的分支條件,改寫為使用數學運算進行無分支化的查詢和計算。例如,用 result = (condition)*a + (!condition)*b 來取代if-else,消除了分支指令本身,但增加了計算量。
  3. 執行緒束級投票與協作:利用__ballot_sync等執行緒束內部通訊原語,讓執行緒束內的一個執行緒先收集所有執行緒的路徑需求,然後驅動整個執行緒束順序執行各個路徑。這與硬體掩碼機制類似,但允許更復雜的協作模式。

通過這些軟硬體結合的演變,SIMT架構正在從一個“必須避免發散”的脆弱模型,進化為一個“能智慧處理發散的彈性模型”,這使得以其為基礎的GPU能夠越來越自信地應對通用、非規則的平行計算挑戰。


8. 硬體微架構實現:從流多處理器到Tensor Core引擎

將SIMT的抽象哲學轉化為矽晶片上具體執行的物理現實,是現代GPU硬體工程學最令人驚歎的成就之一。這涉及到從頂層GPU晶片的整體版面配置,到一個流多處理器內部的精妙電路設計,再到專門加速人工智慧工作負載的張量計算引擎的層層堆疊。

晶片級架構:多級可擴充套件陣列 一個現代AI GPU(如NVIDIA H100)在宏觀上是一個大規模的對稱多處理器系統。其核心是數十甚至上百個獨立的流多處理器(SM),它們通過一個多級的、高速的片上交叉網際網路絡共享一個大型的L2快取分割槽以及多個高頻寬記憶體(HBM)控制器。每個SM都是一個包含完整取指、解碼、發射、執行、訪存功能的自主計算核心。這種架構決策使得晶片規模極易擴充套件:從消費級顯示卡的幾十個SM到資料中心級GPU的上百個SM,軟體無需任何修改,因為求解問題的執行緒網格自動分佈到所有可見的SM上。這種以SM為單位的高度模組化設計,也是製造大晶片時提升良率的有效手段。

流多處理器內部:四分割槽處理與指令併發 深入到單個SM內部,它並非一個由32個簡單ALU組成的平面陣列,而是一個高度結構化的微架構。以NVIDIA Ampere/Ada Lovelace架構為例,一個SM被劃分為四個相同的處理分割槽(Processing Blocks)。每個分割槽擁有自己獨立的:

  • 執行緒束排程器:從該分割槽追蹤的多個活動執行緒束中選擇一個就緒的進行發射。
  • 分發單元:將指令的操作碼和運算元地址分發給該分割槽的功能單元。
  • 暫存器檔案:提供高速存取該分割槽內執行緒的私有暫存器。
  • 執行單元管線:包含一套16位的FP32/INT32 CUDA Core,一套FP64 Core,一套載入/儲存單元,以及特殊功能單元。

這種四分割槽的設計使得一個SM在每個時鐘週期可以同時選擇和發射多達四個不同執行緒束的指令(每個分割槽一個),只要這些指令沒有資料依賴和功能單元衝突。這是指令級並行在SIMT世界的獨特變體。同時,這種將數百執行緒暫存器分塊管理的策略,避免了建置一個統一、超大埠暫存器檔案帶來的功耗和麵積災難。

CUDA Core與特殊功能單元 每個處理分割槽的CUDA Core是SIMT模型最純粹的體現。早期的CUDA Core支援全速FP32和INT32。隨著AI訓練對FP16/INT8的精度需求激增,現代CUDA Core普遍增加了對多精度打包運算的支援,能夠在同一個暫存器(32位)內打包兩個FP16運算元,以雙倍的速率進行吞吐計算。特殊功能單元則專門用於執行超越函式(如sin, cos, exp, rsqrt),這些單元的吞吐量低於普通運算,但在一個執行緒束內所有執行緒都需要這些函式時能提供關鍵加速。

Tensor Core:異構融合的AI發動機 這是近現代AI效能飛躍的核心硬體。Tensor Core不是一個獨立的晶片,而是深度整合在SM內部的一個專用功能單元矩陣。以H100的第四代Tensor Core為例,它針對FP8資料格式(E4M3/E5M2)進行了極致最佳化。從SIMT的角度看,一個執行緒束內的32個執行緒仍然協同工作,但它們不再各自生成自己的乘加結果,而是共同構造並向Tensor Core提交一個完整的、例如16x16x16的矩陣分片任務。Tensor Core內部由數以千計的乘加累加器組成脈動陣列,可以在一個時鐘週期內完成這個16x16x16的矩陣乘法與加法。其吞吐量是同時期標量CUDA Core的數倍到數十倍。值得強調的是,Tensor Core的任務分發、結果寫回、同步機制依然無縫地融合在SM的SIMT排程、暫存器和共享記憶體架構中。它是SIMT這一通用元架構為特定領域計算負載進行的極致硬化最佳化,是“通用性”與“專用性”架構融合的典範。

9. 互聯與分割槽:晶片級、節點級與系統級SIMT擴充套件架構

單個GPU無論多麼強大,也無法處理萬億引數級的模型。因此,SIMT的擴充套件哲學必須從晶片內部走向晶片之間、節點之間,直至形成萬卡級的資料中心叢集。這是一個將數百億個SIMT執行緒跨物理邊界協同工作的系統工程難題。

片內網路:多SM的一致性訪問 在GPU內部,所有SM通過一個高頻寬、低延遲的交叉連線網路(Crossbar或Mesh)與L2快取分割槽和記憶體控制器相連。這個網路的拓撲和頻寬是支撐數千個執行緒束併發訪問全域性記憶體和資料共享的基礎設施,是晶片內部的一個微觀“網際網路”。它必須保證來自不同SM的合併訪存請求能夠被高效、公平地送達目標記憶體通道,並保持快取一致性的單層視角(GPU通常沒有複雜的多核對等快取一致性協議,而是依賴更寬鬆的記憶體模型和顯式同步)。

片間互聯:NVLink與Infinity Fabric 為了在節點內實現多GPU間的超高速資料共享,NVIDIA開發了NVLink。這是一種高頻寬、低延遲、支援快取一致性的片間互聯,使得兩塊或更多GPU可以共享各自的HBM池,形成一個統一的大記憶體空間。對於Simultaneous Multi-Node的大型模型訓練,NVLink將多卡連線成一個超大的虛擬SIMT機器。當模型張量跨卡進行分片時,Tensor Core產生的資料可以通過NVLink在另一塊卡的記憶體中直接彙總。AMD的Infinity Fabric扮演類似角色。這種高速片間直連是避免通訊瓶頸吞噬平行計算增益的關鍵。

節點間互聯:InfiniBand與NVSwitch 在節點以外,系統進一步通過高速網路(如NVIDIA的Quantum InfiniBand或Spectrum Ethernet)連線成叢集。這些網路提供400Gb/s甚至800Gb/s的頻寬,並支援RDMA(遠端直接記憶體訪問),允許一個節點上的GPU直接通過網路將資料寫入另一個節點上GPU的記憶體,無需經過雙方CPU和作業系統的笨重介入。NVIDIA的NVSwitch晶片則是連線多個NVLink通道的無阻塞交換器,在有8個或更多GPU的DGX系統中,NVSwitch允許任何GPU以全頻寬與任何其他GPU通訊,打破了點對點直連的頻寬限制。這一套從片內網路到NVLink、NVSwitch再到InfiniBand的互聯體系,讓演算法設計者能將規模擴充套件到百萬張卡,實踐資料並行、張量並行、流水線並行和專家並行等混合策略,這一切的底層,都是無數個SIMT執行緒在跨越物理界限,協同完成一個規模空前宏大的單一計算任務。


10. 程式設計模型與編譯器:從CUDA C++到隱式並行的抽象階梯

SIMT硬體的巨大潛力,需要一套強大、易用且不斷演進的軟體棧來兌現。從底層的指令集架構(ISA)到頂層的Python DSL,整個軟體體系的目標始終如一:在保持效能的同時,將大規模並行程式設計的複雜性一步步隱藏和自動化。

SASS與PTX:硬體與虛擬機器的雙層ISA NVIDIA GPU內部的真實指令集是SASS(Shader Assembler),它直接控制硬體執行。然而,為前向相容和解耦軟硬體迭代,NVIDIA定義了一個穩定的中間指令集PTX(Parallel Thread Execution)。CUDA編譯器NVCC將C++核心程式碼編譯為PTX碼,隨應用程式釋出。當程式在特定GPU上啟動時,GPU的JIT(即時編譯)驅動程式將PTX程式碼最終編譯為目標GPU的SASS碼。這類似Java的位元組碼機制。PTX顯式暴露了SIMT特性,包括執行緒束、暫存器操作、共享記憶體等,但它仍是虛擬化的,確保為舊架構編寫的程式碼能在新架構上繼續執行並獲得性能最佳化。

CUDA C++:程式設計師的顯式控制台 CUDA C++是控制力最直接的程式設計模型。它要求程式設計師顯式管理主機(CPU)與裝置(GPU)間的記憶體複製、定義核心啟動的執行緒網格維度(grid, block),並在核心程式碼中使用threadIdx.xblockIdx.x等內建變數進行工作空間索引與計算。它還將記憶體空間明確區分開來:__global____device____shared____constant__,讓有經驗的開發者能精細調優記憶體層次的使用。這是一個為效能而設計的“顯式手柄”模型,是建置高效能庫(如cuBLAS, cuDNN)的核心工具。

宣告式與隱式程式設計範式:C++並行策略與架構融合 CUDA C++太重。因此,出現了更上層的抽象。NVIDIA的C++標準庫並行策略(如thrust::transform)允許用類似STL的演算法描述並行操作,由庫在底層處理核心啟動。更進一步,OpenACC和OpenMP的GPU目標解除安裝指令允許程式設計師在現有C/Fortran程式碼上新增#pragma acc parallel loop之類的指令,利用編譯器自動完成資料對映和核心生成。這一層抽象是為主流科學計算和工業程式快速遷移到GPU而設計。

AI架構融合與運算元庫:隱式並行的終極形態 在現代人工智慧開發中,90%以上的開發者觸及的最後一公里是PyTorch、TensorFlow等架構。在這個層級,SIMT和CUDA的蹤跡幾乎完全消失了。開發者呼叫F.linear(x, w, b),架構的排程器會分析張量維度、裝置型別和資料精度,從cuBLAS、cuDNN等高度最佳化的底層運算元庫中呼叫一個匹配的函式。而cuBLAS內部的GEMM實現,才是那個精心設計的記憶體分片、執行緒束級暫存器最佳化和Tensor Core排程的SIMT程式碼集合。這個逐層下放的生態系統,讓硬體的SIMT能力轉化為了純粹的張量操作語義,實現了平行計算的終極民主化。


11. 深度學習的SIMT核心庫與應用生態系統

人工智慧在GPU上的繁榮,其直接推手是一系列高度最佳化的、建置於SIMT程式設計模型之上的線性代數與深度學習基座庫。這些庫是演算法和硬體之間的“中介軟體”,它們封裝了SIMT的複雜度,並向上提供了面向數學概念的、高效能的API。這就是NVIDIA的“CUDA護城河”最核心的、充滿粘性的部分。

cuBLAS:線性代數的基石 cuBLAS實現了一套最基礎的稠密線性代數運算,包括向量、矩陣-向量、矩陣-矩陣運算。它是整個AI計算棧的最底層。一個看似簡單的GEMM呼叫,在cuBLAS內部會根據問題規模、矩陣版面配置和精度,動態選擇數百個精心人工調校的SIMT核心中的一個。這些核心是分片策略、暫存器最佳化、共享記憶體bank衝突避免等技術的集大成者。可以說,整個深度學習世界的速度,直接取決於cuBLAS中GEMM實現的速度。

cuDNN:深度學習的加速引擎 如果說cuBLAS是通用磚塊,cuDNN (CUDA Deep Neural Network) 就是為深度學習專門設計的預製件。它為卷積、池化、歸一化、啟用函式、RNN/LSTM等常見DNN層提供了深度最佳化的實現。cuDNN的代表性特性是自動傅立葉變換卷積、Winograd卷積、以及直接卷積的多種演算法選擇。針對特定的輸入大小、卷積核大小和步幅,cuDNN會在其內建的演算法引擎中進行快速基準測試(自動調優),然後選出延時最低的那個SIMT核心來執行。這使得動態多變的模型結構總能近乎最佳地利用硬體。

NCCL:多卡多節點的通訊支柱 單卡算力有限,大型模型必須跨越卡和節點的邊界。NVIDIA Collective Communications Library (NCCL) 由此誕生。它實現了AllReduceAllGatherReduceScatter等核心集體通訊原語。NCCL的極致之處在於,它精妙地利用了GPU內部的SIMT能力、NVLink/NVSwitch的高頻寬以及InfiniBand的RDMA功能,為大規模並行資料傳送設計了複雜的傳輸和歸約演算法(如Ring演算法和Tree演算法)。訓練一個GPT-4模型時,成千上萬張GPU上的SIMT執行緒,並非只為自己計算,更在NCCL的指揮下,利用NVLink直接在環形拓撲中傳遞和累加彼此的梯度,這是SIMT從計算單元擴充套件為通訊單元的有力證明。

TensorRT:推論的終極武器 在模型部署階段,目標是最大化吞吐量、最小化延遲。TensorRT (NVIDIA TensorRT) 接收一個訓練好的模型,並對計算圖進行一系列的編譯時最佳化:消除無用層、合併卷積/偏置/啟用函式為單一核心(運算元融合)、量化浮點模型至INT8或FP8精度、並基於特定目標GPU架構調優核心、生成一個可部署的最佳化引擎。這一引擎的每一次推論呼叫,背後都是一個極致最佳化的SIMT核心序列。這是SIMT在不同生命週期(從靈活的研發訓練到固定的工業部署)展現適應性的範式。

這些庫的共同作用,使得一個在Python中用三行程式碼定義的神經網路層,能夠被最終翻譯成一個在矽片上數千個核心間精密協同、極高效率執行的SIMT指令序列,構成了從學術創新到工業落地的完整閉環。


12. 與非SIMT架構對比:CPU、脈動陣列與資料流的角力

要客觀評估SIMT,必須將它置於整個人工智慧計算架構的競爭版圖中,與其他主流和新興範式進行系統性比較,分析其各自的比較優勢與結構性侷限。

SIMT vs. 超標量/多核CPU CPU核心追求的是單執行緒的極致低延遲。其大容量私有快取、亂序執行、複雜分支預測和極高主頻的代價是昂貴的面積和功耗開銷,限制了單晶片上能整合的核心數(通常最多128核)。而SIMT GPU的SM則犧牲了單執行緒的舒適度(高延遲、順序執行),將寶貴的電晶體用於堆疊海量的簡單執行單元和大頻寬記憶體控制器,以在給定功耗和晶片面積下,獲得數倍於CPU的總吞吐量。在A乘B這類規則負載上,CPU的延遲最佳化策略無用武之地,而GPU的吞吐最佳化策略完美匹配。然而,在涉及大量序列邏輯依賴、不規律的資料訪問或複雜的作業系統互動時,CPU依然是王者。因此,未來的異構計算中心幾乎總是CPU+GPU耦合。

SIMT vs. 脈動陣列 脈動陣列(如Google TPU v1的核心)是另一個極端。它是一個由高度簡化的處理單元組成的網格,資料以節拍方式在單元間傳遞和計算,專為矩陣乘法而硬化設計。其優點是極高的能效和麵積利用效率,因為控制邏輯極其簡單。其缺點是缺乏靈活性。而在SIMT架構內,Tensor Core本質上就是小型的區域性脈動陣列,但其任務分發和結果寫回完全受龐大的SIMT通用執行緒引擎控制。這使得NVIDIA能夠在一次核心任務中,無縫地交織使用脈動陣列(通過Tensor Core的wmma指令)處理GEMM,用標量CUDA Core處理啟用函式或自定義的複雜元素運算,再用載入/儲存單元搬運資料。這種**“通用的”專用性**,是SIMT架構能快速適應層出不窮的AI新運算元(如SwiGLU啟用、RoPE位置編碼等)的關鍵原因,而純脈動陣列或更專用的資料流AI晶片通常在面臨新運算元時,只能退回到效率低下的向量通用處理來補足。

SIMT vs. 粗粒度可重構架構與資料流 可重構架構和資料流AI晶片代表了另一個方向,試圖在靈活性和效率間找到新平衡。它們高度依賴於編譯器將計算圖直接對映到空間陣列上,最理想的情況是一個運算元完全佔據晶片。但當面對控制流複雜、動態稀疏計算或張量形狀多變的負載時,其資源對映和排程效率會受到挑戰。而SIMT的天才之處在於:通過硬體排程器對成千上萬個執行緒束的快速搶佔和切換,它能以一種天然的、分時多工的方式,平滑地吸收這些動態變化和利用率窪地,保持整體的高吞吐。儘管這可能在絕對能

source: 公開揭露與公開資料整理 本頁僅用於產業鏈學習、資訊檢索和研究輔助;不構成投資建議,不預測漲跌,不提供買賣、部位或目標價建議。
完整概念頁 複盤 13 節結構 公司投研頁 沿產業鏈找到受益公司 投資課 把概念轉成可跟蹤模型