執行緒塊

3秒看懂

執行緒塊（Thread Block / CTA） 是 GPU 並行程式設計模型中最核心的組織單元。它將數百個執行緒編為一組，這些執行緒共享同一塊高速片上儲存器（共享記憶體），能夠在塊內快速同步與協作。在深度學習工作負載中，矩陣乘法、卷積等關鍵運算元正是通過精心劃分執行緒塊，把海量計算對映到上千個核心上執行，從而實現遠超 CPU 的訓練/推論吞吐。

3分鐘產業解釋

GPU 之所以能顛覆深度學習，關鍵在於它把「成千上萬並行執行緒」變成了現實。但直接管理幾萬個執行緒既困難又低效。於是 NVIDIA 在 CUDA 中引入 執行緒塊（Cooperative Thread Array, CTA），將執行緒劃分成一個二維或三維的邏輯塊。一個 kernel 啟動時，會生成一個執行緒塊網格（Grid），每個執行緒塊執行相同的指令，但處理不同的資料塊。塊內執行緒可以通過 __syncthreads() 輕量級同步，並通過共享記憶體交換資料。塊與塊之間則無依賴關係，可獨立排程於流多處理器（SM）上。

產業意義：

• 所有主流深度學習架構（PyTorch、TensorFlow）的底層運算元均由執行緒塊實現，例如 cuBLAS 的 GEMM、cuDNN 的卷積。
• GPU 架構演進（Volta、Ampere、Hopper）中，執行緒塊的資源上限（最大執行緒數、共享記憶體、暫存器）是決定運算元效能的關鍵約束。
• 大規模模型訓練（千卡叢集）的通訊與計算重疊，依賴於對執行緒塊執行順序與資源佔用的精細控制。

換句話說，不瞭解執行緒塊，就很難理解運算元最佳化為何如此複雜，以及為何一張 GPU 的“實際算力”與理論峰值存在差距。

15分鐘專家深入

從專家視角看，執行緒塊遠不止“執行緒的籃子”，它直接制約著 GPU 的 佔用率 (occupancy)、延遲隱藏 和 資料區域性性。

關鍵機制：

1. 排程與 Warp：硬體排程器以 32 個執行緒為一組（一個 warp）作為最小執行單元。一個執行緒塊內包含多個 warp（例如 256 個執行緒即 8 個 warp）。SM 會同時駐留多個執行緒塊，當某個 warp 因訪存延遲或同步停滯時，SM 可瞬時切換到另一個 ready 的 warp，將算術操作的吞吐填補延遲空洞。這是 GPU 隱藏記憶體延遲的核心手段。
2. 資源約束：每個執行緒塊會消耗一定量的暫存器和共享記憶體。SM 的資源總量固定，因此單一執行緒塊申請過多資源，會降低 SM 上可同時駐留的執行緒塊數量，從而降低佔用率，減少延遲隱藏的機會。“用共享記憶體換全域性記憶體訪問”與“保持高佔用率”之間往往存在折衷。
3. 層級並行對映：在深度學習運算元中，常採用三層次分解：任務 → 執行緒塊網格（Grid）→ 單個執行緒塊。例如矩陣乘法，Grid 維度對應輸出矩陣的分塊，每個執行緒塊負責計算一個小輸出塊；塊內執行緒再協同從全域性記憶體載入資料到共享記憶體，然後執行內積迴圈。卷積的 Implicit GEMM 或 Winograd 演算法同樣遵循此分層。
4. 同步代價：__syncthreads() 是塊內 barrier，會強迫 warp 停滯，等待塊內所有執行緒到達該點。若塊內執行緒因分支（如不同 if-else 路徑）到達 barrier 的時間參差，產生顯著的同步開銷。因此高效能 kernel 設計會極力避免執行緒塊內的控制流分歧。

與深度學習負載的深度繫結：

• 矩陣乘法 (GEMM)：執行緒塊處理的輸出子塊常設計為 128×128 或 256×128 等大小，而執行緒塊本身的維度通常為 16×16 或 32×8（匯流排程數 256）等，每個執行緒計算輸出矩陣的一個或幾個元素，通過共享記憶體分塊（tiling）複用資料，減少 HBM 流量。
• FlashAttention：通過將注意力計算拆分為多個執行緒塊，每個塊獨自在片上 SRAM 完成部分 softmax 歸一化，然後合併，避免寫入昂貴的 HBM 中間張量。執行緒塊的精確尺寸與共享記憶體使用量決定了該運算元能否在 40KB~100KB 的有限空間裡完成高效計算。
• Megatron-LM 張量並行：通訊原語（AllReduce/ReduceScatter）通常預先呼叫精心最佳化的 kernel，其執行緒塊引數針對不同矩陣大小和資料版面配置調校，以在計算與通訊之間重疊最優。

因此，一個深度學習系統工程師或編譯器開發者，必須能從運算元計算圖的形狀，推演出合理的 Grid 和 Block 維度，並在暫存器、共享記憶體、佔用率三者間尋找 Pareto 最優點。

技術原理（最深）

硬體執行模型

GPU 裝置包含多個 SM，每個 SM 內部有：

• 整數/浮點執行單元、載入/儲存單元、張量核心（Tensor Core）
• 暫存器檔案（如 65536 個 32 位暫存器）
• 共享記憶體/L1 快取（可配置大小，如 128KB 組合可選）

一個 kernel 啟動時，Host 指定 Grid 維度（gridDim）和 Block 維度（blockDim），生成 gridDim.x × gridDim.y × gridDim.z 個執行緒塊，每個塊內含 blockDim.x × blockDim.y × blockDim.z 個執行緒。

  Grid (二維，處理整個輸出張量)
+---------------------+
|  B0  |  B1  |  B2  |
+---------------------+
|  B3  |  B4  |  B5  |   <-- 每個框是一個執行緒塊
+---------------------+

   Block (二維，例如 16x16 執行緒)
+------------------+
| t0  t1 ... t15 |
| t16 ...     t31|
| ...            |
| t240 ...  t255 |
+------------------+

資源對映與佔用率

SM 的資源空間被劃分為“槽位”(slot)，每個執行緒塊佔據一個槽位，可容納一個或多個 warp。

• 最大駐留執行緒塊數 = min( max_blocks_per_SM, total_registers / (regs_per_thread × block_size), shared_mem_per_SM / shmem_per_block )
• 佔用率 = (活躍 warp 數) / (SM 最大支援 warp 數)

例如，某 SM 可容納 64 warp（2048 執行緒）。若每個執行緒塊含 256 執行緒（8 warp），且 SM 最多可駐留 8 個塊，則最大活躍 warp 數為 64，佔用率 100%。若某個塊使用過多共享記憶體，導致 SM 只能駐留 4 個塊，活躍 warp 數降至 32，佔用率 50%，於是 SM 隱藏延遲的能力腰斬，儘管單塊計算強度很高。

深度學習運算元的執行緒塊設計範式

矩陣乘法 C = A × B（A: M×K, B: K×N）

• Grid 維度：(N / TILE_N, M / TILE_M)
• Block 維度：(TILE_M, TILE_N) 中的執行緒數，通常二維版面配置，如 16×16。
• 每個執行緒塊分步將 A 的子塊和 B 的子塊載入到共享記憶體，塊內所有執行緒協作完成該輸出子塊的計算。
• 使用 float4 向量載入、暫存器預取、雙緩衝等技術進一步提升吞吐。

卷積（Im2col + GEMM 或 Winograd）：

• 執行緒塊的摺疊與迴圈展開極為複雜，常由自動調諧器（如 Triton、TVM）搜尋最佳配置。
• 一個趨勢是使用 Hopper 架構的 Thread Block Cluster，將多個執行緒塊宣告為一個叢集，允許跨塊直接訪問對方的共享記憶體（通過 SM-to-SM 網路），進一步減少全域性記憶體往返，這用於加速注意力機制等運算元。

與 Tensor Core 的耦合

從 Volta 起，張量核心 (Tensor Core) 提供 warp 級矩陣乘累加指令（mma.sync）。執行緒塊需要按 warp 為單位組織資料以喂入張量核心。例如，一個執行緒塊包含 8 個 warp，每個 warp 可能需要多次呼叫 mma 指令（單次 mma 僅完成較小形狀，如 Ampere 的 fp16 規格為 m16n8k16），整個 warp 通過多次協作計算才能覆蓋 16×16×16 等更大形狀，然後協作儲存結果。這裡的同步與資料排布精確到每個 warp 的 lane（執行緒），設計難度極高，但也是達到接近峰值 TFLOPS 的必經之路。

技術演進史

• G80 (2006)：首次引入執行緒塊模型，塊內共享記憶體與 __syncthreads()。可啟動的 Grid 為二維，Block 為三維，最大執行緒塊 512 執行緒，共享記憶體 16KB。基於此實現早期卷積網路加速。
• Fermi (2010)：增加 L1/共享記憶體可配置（16KB/48KB），最大執行緒數提升至 1024，更靈活的並行排程使得大規模的矩陣庫（如 cuBLAS）開始成熟。
• Kepler (2012)：引入動態並行（Dynamic Parallelism），允許 kernel 內部啟動新的 kernel 和 Grid，但執行緒塊模型本身保持不變。多 warp 排程器改進，增強了對不規則計算的適應性。
• Pascal (2016)：統一記憶體、NVLink 助力多 GPU 通訊，但執行緒塊模型無本質變化，XLNet、早期 GPT 的分散式訓練仍基於傳統的 Block-Grid 對映。
• Volta (2017)：推出 Tensor Core，引入了 warp 級協作概念，執行緒塊內需要小心翼翼地將資料版面配置適配到 mma.sync。執行緒排程引入獨立執行緒排程，使每執行緒可擁有自己的程式計數器，但 __syncthreads() 語義不變。共享記憶體增大至最高 96KB。
• Ampere (2020)：第三代 Tensor Core，支援稀疏矩陣加速，共享記憶體頻寬進一步提升。同時引入了 非同步複製（cp.async），允許直接從全域性記憶體載入資料到共享記憶體而無需中間暫存器，大幅緩解了執行緒塊載入流水線的停滯。
• Hopper (2022)：革命性變化——執行緒塊叢集（Thread Block Cluster）。多個執行緒塊可以在硬體層面進行同步（Cluster Group Sync），並相互訪問對方的共享記憶體。這使得跨執行緒塊協同成為可能，顯著降低了如 FlashAttention 這類跨塊歸約運算元的 HBM 流量。此外還有分散式共享記憶體 (DSMEM) 概念，為分散式推論和 MoE 中的 all-to-all 通訊提供底層支援。
• 未來：隨著 Chiplet 和 3D 封裝演進，執行緒塊可能跨越物理 Die，統一的地址空間與更細粒度的同步將開啟新的並行範式，進一步影響大型模型的訓練效率。

技術路線對比

維度	NVIDIA CUDA 執行緒塊 (CTA)	AMD ROCm Work‑Group	OpenCL Work‑Group	備註
組織方式	塊內執行緒可三維索引	同樣支援三維索引	三維索引	基本等價，遷移成本低
記憶體模型	共享記憶體 (__shared__) + warp shuffle	區域性資料共享 (LDS) + DS Permute	local memory + sub‑group shuffle	AMD 的 LDS 與 CUDA 共享記憶體在量級上類似（~64KB 可配置），但部分架構有差異
同步指令	__syncthreads(), __syncwarp()	__builtin_amdgcn_s_barrier	barrier(CLK_LOCAL_MEM_FENCE)	功能相似，都需要開發者顯式呼叫來保證記憶體可見性
資源限制	最大 1024 執行緒/塊，暫存器/共享記憶體限制因架構而異	最大 1024 工作項/組，資源相似	最大工作項規格由裝置決定	無本質差異
跨塊協作	Hopper 支援 Thread Block Cluster	當前架構無原生跨 work‑group 的共享記憶體訪問	不支援	NVIDIA 暫時領先叢集式協作
深度學習加速	與 Tensor Core 深度耦合，cuDNN/cuBLAS 極致最佳化	通過 MIOpen/rocBLAS 對映到 Matrix Core，生態略遜但快速追趕	無專用 AI 庫，多作為底層後端	NVIDIA 軟體棧成熟度領先

另類對比：CPU 的並行任務劃分（如 OpenMP 並行 for）不提供等價於執行緒塊內共享記憶體的顯式抽象，因而資料複用靠快取自動管理，無法精細控制，難以達到 GPU 級別的能效比。TPU 等專用架構則採用更粗粒度的脈動陣列，其控制單元對標執行緒塊網格，但內部可程式設計性弱於 CUDA 執行緒塊。

上下游

• 上游：
  • 編譯器與工具鏈：NVCC 編譯器負責將 CUDA 程式碼中的 <<<grid, block 啟動語法翻譯為驅動可以理解的 ABI 呼叫，並分配暫存器和共享記憶體。PTX（Parallel Thread Execution）作為中間表示，向上承接高層語言的執行緒塊抽象，向下對映到 SM 微架構。
  • 運算元庫與架構：cuBLAS、cuDNN、CUTLASS 等庫內部高度模板化，通過 C++ 模板引數在編譯期確定 Block 維度、Tile 大小等，供 PyTorch、JAX 呼叫。新一代 DSL（如 Triton）允許使用者在類似 NumPy 的語義上定義定址與分塊，編譯器自動生成執行緒塊配置，極大降低開發門檻。
• 下游：
  • 大型模型訓練（GPT‑4、Llama 3 等）：所有混合精度訓練、TP/PP/DP 拆分後的矩陣運算，最終都分解為無數個執行緒塊的執行。執行緒塊的資源利用率直接影響單卡訓練成本。
  • 推論服務：batch 推論中的 GEMM 與 Attention 運算元，需要低延遲最佳化，往往使用持久化執行緒塊（Persistent Thread Block）技術避免 kernel 重複啟動開銷。
  • 邊緣與自動駕駛：低功耗 GPU（如 Orin）上的執行緒塊配置需極度節省暫存器，以在有限的 SM 頻寬上跑即時模型。

關鍵指標

• 執行緒塊尺寸：典型值為 128、256、512 個執行緒，二維版面配置（如 16×16 或 32×8）常用在矩陣運算中。過小則 warp 浪費，過大則資源不足。
• 共享記憶體使用量 (per block)：直接影響佔用率。深度學習 kernel 通常在 16KB~48KB 範圍內取平衡。若需超過硬體上限，必須退化為多次小分塊或使用全域性記憶體。
• 暫存器壓佔 (register pressure)：編譯器分配給每個執行緒的暫存器數量。例如，每個執行緒使用 128 個暫存器，塊大小 256，則消耗 256×128 = 32768 暫存器，可能成為 SM 容納塊數的瓶頸。
• 佔用率：工程師常通過 NVIDIA Nsight Compute 工具檢視 “Occupancy” 百分比，理想值 >50%，但低佔用率且高算數強度時也可能仍達到高效（如某些 GEMM kernel 僅需 2-3 塊即可填滿計算單元）。
• warp 排程效率：由於 warp 是排程單元，塊內 warp 數量最好是 2 的冪且≥ 4（128 執行緒），保證排程器有足夠的“待命 warp”。
• 同步開銷：塊內 barrier 使用頻繁或分歧嚴重，將導致 SM 空閒週期。無硬數字，需通過 profiler 測量 stall 原因。

供需與市場資料

由於檢索失敗，無法提供具體的 GPU 出貨量或利用率資料。但可以定性描述：

• GPU 供應直接影響執行緒塊並行規模：A100/H100 等資料中心 GPU 持續供不應求，每張卡的 SM 數量、共享記憶體大小決定了可以同時駐留的執行緒塊總量。全球 AI 算力需求劇增的情況下，雲端廠商與科技巨頭對 Hopper 架構 GPU 的爭奪，實質上是在爭奪可併發執行緒塊的數量與訪問頻寬。
• 架構代際供給：NVIDIA 的安培→霍普→布萊克威爾的迭代，每次帶來 ~1.5‑2.7 倍的執行緒塊硬體資源提升（如 H100 的叢集塊、更大共享記憶體頻寬）。這直接縮短大型模型單步訓練時間，降低總擁有成本（TCO）。
• 軟體生態的“執行緒塊效率紅利”：架構代際更新往往需要運算元重新調優（例如 FlashAttention-2、CUTLASS 3.x 專門為 Hopper 的叢集塊設計），因此能快速適配新架構的架構和團隊將獲得顯著的算力利用優勢，推動頭部 AI 公司的資本回報。

代表公司與資本對映

• NVIDIA：CUDA 執行緒塊模型的建立者和持續推動者。從 G80 到 H200，執行緒塊實現細節的逐代進化構築了其強大的軟體護城河。資本市場上，輝達資料中心營收的爆發式增長，直接受益於其 GPU 上高效執行緒塊執行帶來的深度學習訓練成本優勢。
• AMD：通過 ROCm 生態追趕，其 work‑group 模型與 CUDA 執行緒塊高度相容。Instinct MI300X 等產品若能在軟體棧層面彌合 CUTLASS/hipBLAS 的效能差距，有望從 AI 算力市場中分得份額。
• 新興 GPU/加速器公司（摩爾線程、壁仞等）：均需實現一套等價於執行緒塊的並行機制。能否提供與之匹配的編譯器及標準庫，決定了市場對其“國產替代”效能的認可。
• 雲端廠商：AWS Trainium、Google TPU 等專用晶片雖然不直接暴露執行緒塊概念，但其底層多維並行單元的設計參考了相似的分層對映思想。資本關注點在於它們能否以更低成本承載原先由 GPU 執行緒塊承擔的注意力/矩陣運算負載。

產業驗證邏輯

• 執行緒塊效率驗證：GPU 算力不會自動轉化為訓練吞吐，取決於運算元庫對執行緒塊資源的使用效率。評估深耕 CUDA 生態、能快速推出適配新架構庫的團隊時，應重點看其 kernel 適配速度、吞吐提升幅度和客戶驗證結果。
• 關注軟硬協同：每次 GPU 架構引入新的執行緒塊特性（如 Hopper 的 Tensor Memory Accelerator、Clusters），都要求軟體重寫關鍵 kernel。能夠快速完成這一過程的公司，將在同期硬體採購中獲得先發成本優勢。
• 風險點：若未來演算法趨勢（如 MoE、動態路由）導致計算模式朝著更細粒度、不規則的方向發展，傳統密排執行緒塊的優勢可能減弱，部分專用稀疏加速器也可能侵蝕 GPU 在特定推論負載中的份額，需持續追蹤。
• 估值錨點：算力即吞吐，而執行緒塊是吞吐的基本計費單元。評估一個 AI 資料中心的 CAPEX 效率，最終會落到每個 SM 每時鐘執行的執行緒塊數量上。任何導致該指標惡化的架構/軟體缺陷都會反映在運維成本中。

常見誤讀糾偏

誤讀1：「執行緒塊數量越多，並行度越高，效能越好」 正解：效能取決於有效隱藏延遲的足夠多 warp，而非執行緒塊數量本身。若每塊分配的暫存器或共享記憶體過多，SM 只駐留 1-2 個塊，雖然塊數少，但每塊極高效地利用 Tensor Core，仍可能跑滿峰值；相反，大量小塊可能引起暫存器溢位或排程開銷，導致效能下降。最佳化核心本質上是在資源限制內尋找“塊數×每塊效率”的最大積。

誤讀2：「執行緒塊之間可以通過全域性記憶體隨意通訊」 正解：執行緒塊之間沒有執行順序保證，不能使用 __syncthreads() 同步；它們雖然可以通過全域性記憶體通訊，但缺乏高效同步原語（Hopper 之前）。如果在需要跨塊一致性的場景下（如全域性歸約）隨意操作，可能導致資料競爭和不確定結果。正確做法是單個 kernel 內不進行跨塊同步，要麼拆成多個 kernel，要麼使用 Hopper 的叢集塊或協作組（Cooperative Groups）的網格同步。

誤讀3：「所有執行緒塊使用相同的共享記憶體大小」 正解：在同一個 kernel 內，所有執行緒塊確實使用相同尺寸的靜態共享記憶體（編譯期確定）或最大動態共享記憶體（執行時指定）。但是沒有任何機制強制它們實際使用量一致，開發者有責任不越界。如果某個塊內使用超過宣告大小，會導致未定義行為，極難除錯。

學習路徑

1. 基礎：閱讀《CUDA C++ Programming Guide》的“Execution Model”和“Memory Hierarchy”章節，理解 Grid、Block、Thread 的索引計算及共享記憶體宣告。
2. 動手：手寫一個樸素的矩陣乘法（naive GEMM），然後逐步新增共享記憶體分塊、暫存器預取、向量化載入，使用 Nsight Compute 觀察執行緒塊佔用率和warp stall 原因。
3. 進階：研讀 CUTLASS 的 gemm_device_level 和 gemm_block_level 抽象，分析如何通過模板引數靜態配置執行緒塊形狀和計算過程。推薦《CUTLASS: Fast Linear Algebra in CUDA C++》論文。
4. 現代實踐：學習 Triton 語言，通過 triton.autotune 探索執行緒塊配置的自動搜尋，理解為何編譯器能產生接近手寫 CUDA 效能的程式碼。
5. 前沿：深入 Hopper TMA 與 Thread Block Cluster 示例（如 NVIDIA 官方 cuda-samples 中的 cudaTensorCoreGemm 和 FlashAttention-2 原始碼），掌握跨塊直接共享記憶體的程式設計模型。
6. 系統思考：結合 Megatron-LM 或 vLLM 原始碼，理解大型模型系統中的 kernel 呼叫圖譜，建立“執行緒塊 → 運算元 → 分散式訓練”的全域性觀。

一句話總結

執行緒塊是 GPU 上最小但最關鍵的執行協作單元，它的尺寸、資源分配與同步方式直接劃定了深度學習計算的效率邊界；掌握執行緒塊，就掌握了挖掘 AI 算力的鑰匙。

延伸閱讀與來源

• NVIDIA CUDA C Programming Guide (Programming Model) – 官方定義與資源約束
• NVIDIA Hopper Architecture Whitepaper – Thread Block Cluster、DSMEM 等新特性
• CUTLASS Documentation & Repository (github.com/nvidia/cutlass) – 矩陣乘法的執行緒塊分解範式
• FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning – 跨塊/跨 warp 負載均衡的實現
• 各代 GPU 架構微手冊（NVIDIA Fermi/Kepler/Pascal/Volta/Ampere Tuning Guides） – 佔用率與資源調優數值
• 《Programming Massively Parallel Processors》 (David Kirk & Wen-mei Hwu) – 執行緒塊設計案例與習題
• Nsight Compute Kernel Profiling Guide – 如何測量執行緒塊的實際效能瓶頸

（注：因聯網檢索失敗，具體產品資料（如某架構的共享記憶體精確大小）未引出處；文中所述均為行業共識性概念，不構成投資建議。）