CUDA(統一運算架構)
/ KOO-duh /
一旦你有了 GPU 並想在它上面跑你自己的運算,你就需要一種方式來撰寫那個運算、把它送到裝置上,並啟動它的數千份副本。CUDA 就是 NVIDIA 為了在 NVIDIA GPU 上正好做這件事所推出的平台與程式設計模型。你用一種 C/C++ 方言寫一個叫做核心(kernel)的小函式,標記它在裝置上執行,然後請 GPU 把它跨越一整個執行緒網格來執行。
這個模型有一套值得弄清楚的特定詞彙。核心是每個執行緒所執行的函式。你以類似 myKernel<<<numBlocks, threadsPerBlock>>>(...) 的配置來啟動它,這會建立一個由區塊(block)組成的網格(grid),每個區塊又由執行緒組成;每個執行緒都得知自己的座標(blockIdx、threadIdx),並用它來挑出自己負責哪個資料元素——第 i 號執行緒處理 c[i]。在底層,硬體並非一次排程一個執行緒:它以 32 個為一組、叫做一個 warp 的方式讓它們步調一致地執行,這 32 個執行緒一起執行同一條指令。同一區塊內的執行緒可以透過快速的晶片內共享記憶體合作,並用屏障(__syncthreads())同步,而不同區塊彼此獨立執行,甚至可能不會同時存活。主機(CPU)端的程式碼用 cudaMalloc() 配置裝置記憶體、用 cudaMemcpy() 把輸入複製過去、啟動核心,再把結果複製回來。
它之所以重要,是因為 CUDA 比任何其他東西都更使通用 GPU 運算成為主流,至今仍是深度學習與高效能運算的主導生態系,上面疊著一整套深厚的函式庫。要老實談兩件事。第一,CUDA 是專有的、只限 NVIDIA;OpenCL 與 SYCL 這類可攜替代方案的存在,正是因為這種綁定。第二,取得真正的速度並非自動發生:你得避免 warp 在分支上發散、把記憶體佈局成讓一個 warp 的存取合併(coalesce)成少數幾次傳輸、並把資料傳輸與運算重疊——「我寫了一個核心」和「它很快」是兩個非常不同的里程碑。
__global__ void add(const float *a, const float *b, float *c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; // 這個執行緒的索引 if (i < n) c[i] = a[i] + b[i]; // 每個執行緒處理一個元素 } // 主機端:add<<<(n+255)/256, 256>>>(a, b, c, n); // 由每 256 執行緒一區塊組成的網格
一個 CUDA 核心:數千個執行緒各自算出自己的索引、各處理一個陣列元素。主機端啟動一個由區塊組成的網格。
CUDA 是專有的,只能在 NVIDIA 硬體上執行——這正是可攜選項(OpenCL、SYCL)存在的原因。而且寫出一個能編譯的核心很容易;讓它變快卻需要避免 warp 發散、合併記憶體存取,並隱藏 CPU 與 GPU 之間的傳輸成本——這些編譯器都不會替你做。