GPU 與異質計算

為什麼要再來一顆處理器——而且是非常不同的一顆

這整條階梯裡的一切，從儲存程式模型到虛擬記憶體再到執行緒，都悄悄假設了一種引擎：一顆為了把單一指令流跑得快而打造的 CPU，深層的分支預測器、大快取與亂序執行，全都花在讓單一執行緒更早完成上。GPU 押的是相反的賭注。它不靠幾顆聰明的核心去追逐低延遲，而是塞進數千條簡單的算術通道，把一切都押在吞吐量上：對一大片資料一次做同一個運算。CPU 是載一位乘客的跑車；GPU 是一列貨運火車。GPU 計算就是把那列貨運火車拿來做一般計算的實作，而不只是畫像素。

把這類截然不同的引擎混在一起的機器，叫做異質的（heterogeneous）：一個程式部分跑在 CPU 上，部分跑在一個或多個加速器上，而這些加速器擅長 CPU 不擅長的事。這已不再罕見——你的手機、你的筆電、雲端每一個 GPU 節點都是異質的。它對系統程式設計的後果尖銳，也是這整篇導引的主題：你不能再寫單一條指令流、跨單一個位址空間，就把它叫做一個程式。你必須決定哪一個引擎跑哪一部分工作，並在那兩個引擎不共享的記憶體之間搬移資料。這項技能不再只是貼近底層——而是同時貼近兩種非常不同的底層。

SIMT：數千條執行緒如何跑同一個程式

你在 CPU 上已經見過 SIMD：一條指令把同一個運算套用到一小段向量的值上——比方說一個暫存器裡的八個 float。GPU 把這個想法推到極致，並用不同的外衣包裝，這個模型叫做 SIMT，單指令多執行緒。你寫一個小函式——一個核心（kernel）——它看起來像在計算單一個元素：它讀自己的索引、載入一個值、做一些數學、寫出一個結果。GPU 接著把那同一個核心一次跨數千條通道發射，每條通道有自己的索引。SIMD 要你以向量去思考，SIMT 卻讓你寫看似純量的程式碼，再讓硬體跑一大群副本。底下是同一種資料平行性，只是換了一張和善的臉。

// CPU view: a loop over N elements
for (int i = 0; i < N; i++)
    c[i] = a[i] + b[i];

// GPU (SIMT) view: ONE kernel, launched N times,
// each instance handed its own index i
__global__ void add(const float *a, const float *b,
                    float *c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N)            // guard: last block may overhang
        c[i] = a[i] + b[i];
}
// the loop has vanished; the hardware *is* the loop

同一個向量加法，兩種寫法。在 CPU 上迴圈是明寫的；GPU 核心計算單一個元素，由執行期發射 N 份副本，每一份有不同的索引。那個 if 守衛很重要，因為發射數會被無條件進位到整數個區塊，於是有些通道沒有元素可碰。

這裡有一個和善外衣藏起來的誠實陷阱。那些通道並非真正獨立：硬體把它們以鎖步群組的方式跑（NVIDIA 把 32 條一組稱為一個 warp）。群組裡的每一條通道都必須在同一時刻執行同一條指令。所以當你的核心碰上一個 if，有些通道走分支、有些不走時，硬體無法同時跑兩邊——它跑走分支的那些通道，讓其他的閒置，再跑另一邊。這叫做分支發散（branch divergence），是你在 CPU 上學過的分支代價的 GPU 版本，只不過罰則是浪費掉的通道，而非被清空的管線。GPU 上滿是分支的程式碼會悄悄扔掉你大部分的吞吐量，這正是為什麼 GPU 偏愛扁平、一致、算術繁重的迴圈，而厭惡決策樹。

卸載模型：資料必須越過一道鴻溝

GPU 不會跑你的整個程式——它跑 CPU 請它跑的那些核心。CPU 仍然是老闆：它擺好資料、把一個核心與它的引數交給 GPU、等結果、再繼續。這就是卸載模型，而它最難的一個事實，是 GPU 通常有自己的記憶體。一顆獨立 GPU 隔著一條匯流排（PCIe，或封裝上更快的連結）坐著，配著數 GiB 自己的 RAM，而你在 CPU 上用 malloc() 配置的那個位址空間，在那一邊毫無意義。CPU 指標不是 GPU 指標。在一個核心能碰你的陣列之前，那些位元組必須實體地越過匯流排被複製進 GPU 記憶體；它跑完之後，結果又必須複製回來。

在 GPU 記憶體（GPU 自己的 RAM）裡配置緩衝區，與你在 CPU 上 malloc() 出來的陣列分開。
把輸入位元組由主機端向裝置端越過匯流排複製過去——由 CPU 發起，但一個 DMA 引擎搬運那些位元組，不需要 CPU 一個一個照看。
發射核心，指定要產生多少條通道（執行緒，編組成區塊）；這個呼叫立刻返回，GPU 非同步地跑。
同步——等到 GPU 回報核心完成——然後把結果由裝置端向主機端複製回來，到這時才在 CPU 上檢查它們。

那道匯流排的跨越，正是異質效能生死所繫之處。GPU 能在眨眼間把一百萬個 float 加起來，但如果你把那一百萬個 float 複製過去、加完、再複製回來，那些複製會讓計算相形見絀，整件事比 CPU 原本要做的還慢。一條值得背下來的殘酷經驗法則是：卸載只有在每位元組的算術量高到足以遮蓋傳輸時，或在資料能跨許多核心常駐在 GPU 上、讓你複製一次計算多次時，才划算。這跟效能那一級的受計算拘束對受記憶體拘束思考是同一回事，如今被拉伸跨過一條實體匯流排。現代系統用統一記憶體（unified memory）把它模糊化，由執行期在單一個指標背後依需求遷移分頁——但一個分頁首次被觸碰時仍會實體地搬動，所以代價沒有消失，只是被藏起來了。

程式設計模型：CUDA、SYCL 與計算著色器

你究竟怎麼寫一個核心？主流的答案是 CUDA，NVIDIA 的工具組（首次發行於 2007 年），它用幾個關鍵字擴充 C++：__global__ 把一個函式標記為核心，角括號發射語法 kernel<<<blocks, threads>>>(args) 把它產生出來，再由執行期呼叫處理配置—複製—發射—複製的那套舞步。CUDA 成熟又快，但它把你綁在 NVIDIA 硬體上——這是它誠實的代價。可攜的替代方案以一些尖峰效能換取在不只一家廠商的晶片上執行。SYCL 是一個 Khronos 標準，用純粹的、標準的 C++（不擴充語言）表達同樣的卸載，於是同一份原始碼能透過不同的後端瞄準 NVIDIA、AMD 與 Intel 的 GPU。而計算著色器（compute shaders）——在 Vulkan、Metal 或 WebGPU 裡——讓你透過一個繪圖 API 跑一般的 SIMT 核心，這正是前一篇導引裡瀏覽器從 WebAssembly 抵達 GPU 的途徑。

在語法底下，這三者共享你如今已經見過的架構：一個主機程式、為裝置編譯的核心、明寫或被遷移的資料搬移，以及 SIMT 執行。它們甚至共享一段編譯器後端的故事——一個 CUDA 核心會被編譯成一種叫做 PTX 的中間形式，再由驅動程式在載入時定稿成實際 GPU 的指令，很像一種可攜的位元組碼。所以你為 CPU 內化的那座抽象堆疊在這裡重演了一遍，只是高了一層：原始碼、到一個可攜的 IR、到廠商機器碼，只是換了一台機器。真正新的東西不是編譯器管線；而是那個目標是一台隔著匯流排、有上千條通道的吞吐量引擎，而你身為系統程式設計師的工作，是餵飽它而不讓它在資料傳輸上挨餓。

這在前沿裡的位置

退一步，看看為什麼 GPU 是系統前沿的頭條。CPU 大約在二十年前就不再每執行緒地大幅變快了——時脈停滯，而指令級平行性那頓免費午餐也撞上了報酬遞減。業界的答案是走向寬與走向專門：更多核心，接著為如今主導的工作負載打造整個不同的引擎——繪圖、模擬，尤其是位於機器學習核心的稠密線性代數。一次現代的訓練，底下其實就是上面你看到的那種向量加法核心的數十億次，經過融合與分塊，常駐在 GPU 記憶體裡以躲開匯流排。理解卸載模型與 SIMT 已不再是冷門；它正是世界上一大部分計算實際運行的方式。

而 GPU 只是最顯眼的加速器。同樣的異質模式——一個主機加上一個隔著匯流排、配著自己的記憶體與自己指令流的專門裝置——也涵蓋張量單元（TPU）、DSP、FPGA，以及核心旁路網路那篇導引點到的智慧網卡。這個系統課題可以一般化：未來不是一台更快的引擎，而是許多台專門的引擎，而那些困難又持久的技能，是決定什麼跑在哪裡、管理跨不共享記憶體的資料搬移，以及對同時運行的引擎之間的順序與同步進行推理。這條階梯起步時那個單一 CPU、單一位址空間的世界，一直都是一種簡化；異質計算正是那個簡化終於、誠實地畫下句點之處。