1. 作業(yè)背景與核心挑戰(zhàn)從理論到CUDA實(shí)戰(zhàn)的跨越又到了高性能計(jì)算編程的作業(yè)季這次是第五次作業(yè)。如果你和我一樣之前幾個(gè)作業(yè)可能還在用OpenMP、MPI折騰CPU上的并行那么這次作業(yè)大概率會(huì)是一個(gè)分水嶺——我們要真正開始接觸GPU編程也就是CUDA。從網(wǎng)絡(luò)上的熱詞也能看出來“CUDA安裝”、“線程布局”、“shared memory”這些詞被反復(fù)提及這恰恰說明了從理論學(xué)習(xí)轉(zhuǎn)向?qū)嶋H編碼時(shí)大家普遍會(huì)遇到的那道坎環(huán)境配置和概念落地。這份作業(yè)的核心絕不是讓你寫一個(gè)能“跑起來”的Hello World。它的深層價(jià)值在于逼迫你理解GPU這個(gè)眾核怪獸的思維方式。CPU編程是串行思維加點(diǎn)并行優(yōu)化而CUDA編程要求你從一開始就用并行的視角去設(shè)計(jì)數(shù)據(jù)、劃分任務(wù)。作業(yè)里提到的“線程布局”和“內(nèi)存層次”就是CUDA編程的兩大基石。線程布局決定了你的計(jì)算任務(wù)如何被成千上萬個(gè)微小的計(jì)算單元線程消化內(nèi)存層次則決定了數(shù)據(jù)如何高效地在芯片內(nèi)流動(dòng)避免讓高速的計(jì)算核心餓著肚子等數(shù)據(jù)。shared memory共享內(nèi)存作為其中最關(guān)鍵的一環(huán)用得好能讓程序性能飛升用不好或者不用可能比CPU版本還慢。所以面對這個(gè)作業(yè)我們首先要調(diào)整心態(tài)這不是一次簡單的編程練習(xí)而是一次計(jì)算機(jī)體系結(jié)構(gòu)思想和并行編程范式的實(shí)戰(zhàn)訓(xùn)練。目標(biāo)不是交差而是真正弄明白為什么我的矩陣乘法GPU版本在某些情況下可能還不如用OpenBLAS優(yōu)化的CPU版本快問題往往就出在對線程和內(nèi)存的理解深度上。2. 環(huán)境搭建避開“No kernel image”與版本兼容的深坑動(dòng)手寫代碼之前環(huán)境是第一個(gè)攔路虎。熱搜詞里“cuda安裝失敗”、“no kernel image is available for execution”高居前列這幾乎是每個(gè)CUDA新手的必經(jīng)之痛。這個(gè)問題說白了就是你編譯的CUDA代碼內(nèi)核與當(dāng)前GPU的硬件架構(gòu)不兼容。GPU和CPU不同它有所謂的“計(jì)算能力”Compute Capability比如RTX 4060是8.9Tesla P40是6.1。你用為計(jì)算能力8.9編譯的內(nèi)核去一塊計(jì)算能力6.1的老卡上跑就會(huì)觸發(fā)這個(gè)錯(cuò)誤。2.1 精準(zhǔn)確定環(huán)境配置鏈條解決這個(gè)問題需要一個(gè)清晰的配置鏈條GPU硬件 - NVIDIA驅(qū)動(dòng) - CUDA Toolkit - 深度學(xué)習(xí)框架如PyTorch。鏈條中任何一環(huán)版本不匹配都可能導(dǎo)致災(zāi)難。首先用nvidia-smi命令查看你的驅(qū)動(dòng)版本和GPU型號。這個(gè)命令輸出的右上角會(huì)顯示“CUDA Version: 12.4”之類的信息注意這個(gè)不是你安裝的CUDA Toolkit版本而是此驅(qū)動(dòng)最高支持的CUDA運(yùn)行時(shí)版本。你的CUDA Toolkit版本必須等于或低于這個(gè)值。然后去NVIDIA官網(wǎng)根據(jù)你的操作系統(tǒng)和驅(qū)動(dòng)版本選擇對應(yīng)的CUDA Toolkit。對于作業(yè)編程通常選擇最新的穩(wěn)定版如CUDA 12.x即可因?yàn)樗嫒菪宰詈蒙鐓^(qū)支持也最廣。下載時(shí)建議選擇runfile本地安裝方式因?yàn)樗试S你更靈活地選擇安裝組件尤其是在系統(tǒng)已存在多個(gè)CUDA版本時(shí)。2.2 安裝實(shí)操與多版本管理在Linux包括WSL2下安裝步驟大致如下# 1. 賦予安裝文件執(zhí)行權(quán)限 chmod x cuda_12.4.0_550.54.14_linux.run # 2. 運(yùn)行安裝程序關(guān)鍵一步是取消驅(qū)動(dòng)安裝除非你需要更新驅(qū)動(dòng) sudo ./cuda_12.4.0_550.54.14_linux.run安裝界面中你會(huì)看到一堆組件選項(xiàng)。如果你已經(jīng)安裝了合適的NVIDIA驅(qū)動(dòng)務(wù)必取消勾選“Driver”只安裝CUDA Toolkit本身。否則可能會(huì)覆蓋現(xiàn)有驅(qū)動(dòng)引發(fā)顯示問題。安裝完成后需要配置環(huán)境變量。我個(gè)人的習(xí)慣是在~/.bashrc或~/.zshrc中這樣設(shè)置export PATH/usr/local/cuda-12.4/bin${PATH::${PATH}} export LD_LIBRARY_PATH/usr/local/cuda-12.4/lib64${LD_LIBRARY_PATH::${LD_LIBRARY_PATH}}這里有一個(gè)關(guān)鍵技巧不要將/usr/local/cuda這個(gè)軟鏈接路徑放入環(huán)境變量。而是明確指定具體版本路徑如cuda-12.4。這樣當(dāng)你需要切換版本時(shí)只需修改環(huán)境變量指向另一個(gè)具體路徑如cuda-11.8或者切換這個(gè)軟鏈接的指向非常清晰避免了版本混亂。2.3 驗(yàn)證安裝與編譯測試安裝后通過nvcc --version查看編譯器版本用nvidia-smi再次確認(rèn)驅(qū)動(dòng)。然后編譯一個(gè)簡單的測試程序// test_cuda.cu #include stdio.h __global__ void helloFromGPU() { printf(Hello World from GPU thread %d!\n, threadIdx.x); } int main() { helloFromGPU1, 5(); cudaDeviceSynchronize(); return 0; }使用nvcc test_cuda.cu -o test_cuda編譯并運(yùn)行。如果成功打印說明基礎(chǔ)環(huán)境OK。注意如果你在WSL2中操作務(wù)必確保已安裝WSL2專用的NVIDIA驅(qū)動(dòng)并在Windows主機(jī)和WSL2中保持驅(qū)動(dòng)版本大致匹配。WSL2下的CUDA安裝包也需要從NVIDIA官網(wǎng)的WSL2專區(qū)下載。3. 核心概念拆解線程層次、內(nèi)存模型與性能要害環(huán)境搞定后我們來啃硬骨頭CUDA的編程模型。很多同學(xué)看了書上的示意圖覺得線程網(wǎng)格Grid、線程塊Block、線程Thread三層結(jié)構(gòu)很簡單但一到自己設(shè)計(jì)時(shí)就會(huì)懵。關(guān)鍵在于要把這個(gè)抽象模型和你具體的計(jì)算任務(wù)比如矩陣乘法、圖像卷積的數(shù)據(jù)結(jié)構(gòu)結(jié)合起來思考。3.1 線程布局設(shè)計(jì)從數(shù)據(jù)維度出發(fā)CUDA的線程組織是分層且多維的。一個(gè)內(nèi)核Kernel啟動(dòng)時(shí)你指定一個(gè)網(wǎng)格Grid網(wǎng)格由多個(gè)線程塊Block組成每個(gè)塊又包含多個(gè)線程。它們都可以是一維、二維或三維的。設(shè)計(jì)線程布局的第一原則是讓一個(gè)線程處理一個(gè)數(shù)據(jù)元素或一小部分。例如對于一個(gè)MxN的矩陣加法我們可以啟動(dòng)一個(gè)MxN的二維線程網(wǎng)格讓線程(i,j)去處理矩陣C[i][j] A[i][j] B[i][j]。但網(wǎng)格維度有上限如65535 x 65535 x 65535塊內(nèi)的線程數(shù)也有上限通常是1024。所以對于超大矩陣我們需要讓一個(gè)線程處理多個(gè)數(shù)據(jù)。更常見的做法是啟動(dòng)的線程總數(shù)略多于數(shù)據(jù)總數(shù)通過線程ID來映射數(shù)據(jù)索引。例如處理N個(gè)元素我們啟動(dòng)(N255)/256個(gè)塊每個(gè)塊256個(gè)線程。在線程中int idx blockIdx.x * blockDim.x threadIdx.x; if (idx N) { // 處理data[idx] }這里blockDim.x是塊的大小256blockIdx.x是塊的索引threadIdx.x是線程在塊內(nèi)的索引。idx就是全局線程ID我們用它作為數(shù)據(jù)索引。3.2 內(nèi)存層次詳解帶寬與延遲的博弈這是CUDA性能優(yōu)化的核心。GPU內(nèi)存分為多個(gè)層次速度、大小和用法天差地別。全局內(nèi)存Global Memory容量最大GB級別速度最慢延遲最高。所有線程都能讀寫是主機(jī)CPU與設(shè)備GPU數(shù)據(jù)傳輸?shù)闹饕獦蛄?。訪問全局內(nèi)存要盡量合并Coalesced即連續(xù)的線程訪問連續(xù)的內(nèi)存地址這樣硬件可以一次事務(wù)讀取一大塊數(shù)據(jù)極大提升帶寬利用率。共享內(nèi)存Shared Memory位于每個(gè)流多處理器SM片上速度比全局內(nèi)存快數(shù)十倍但容量很小通常每塊幾十KB。同一個(gè)線程塊內(nèi)的所有線程共享這片內(nèi)存。它是手動(dòng)管理的緩存用于存儲線程塊需要反復(fù)訪問的數(shù)據(jù)。例如在矩陣乘法中將矩陣的子塊從全局內(nèi)存加載到共享內(nèi)存然后所有線程從共享內(nèi)存中快速讀取數(shù)據(jù)進(jìn)行計(jì)算能極大減少對全局內(nèi)存的訪問。寄存器Registers速度最快每個(gè)線程私有。用于存儲局部變量。寄存器資源有限如果線程使用的寄存器過多會(huì)導(dǎo)致活躍線程數(shù)減少影響并行度。常量內(nèi)存Constant Memory和紋理內(nèi)存Texture Memory用于特殊訪問模式有緩存機(jī)制。3.3 Shared Memory實(shí)戰(zhàn)以矩陣乘法為例我們以最經(jīng)典的平鋪Tiled矩陣乘法為例看shared memory如何發(fā)揮作用。假設(shè)計(jì)算C A * BA是MxKB是KxN。 沒有優(yōu)化時(shí)每個(gè)線程計(jì)算C的一個(gè)元素需要讀取A的一整行和B的一整列導(dǎo)致對全局內(nèi)存的訪問次數(shù)是O(MNK)且訪問不連續(xù)。采用平鋪優(yōu)化后我們將矩陣分塊Tile。假設(shè)塊大小為TILE_WIDTH如16。那么每個(gè)線程塊負(fù)責(zé)計(jì)算C中一個(gè)TILE_WIDTH x TILE_WIDTH的子矩陣。為了計(jì)算這個(gè)子矩陣需要A中對應(yīng)的一個(gè)行塊和B中對應(yīng)的一個(gè)列塊。我們將這些行塊和列塊從全局內(nèi)存加載到共享內(nèi)存數(shù)組ds_A和ds_B中。同一個(gè)線程塊內(nèi)的所有線程協(xié)同完成加載工作每個(gè)線程加載一個(gè)元素到ds_A和ds_B。然后所有線程同步__syncthreads()確保共享內(nèi)存數(shù)據(jù)加載完畢。接著線程使用共享內(nèi)存中的數(shù)據(jù)進(jìn)行局部乘加計(jì)算。移動(dòng)“平鋪窗口”重復(fù)加載、同步、計(jì)算的過程直到處理完所有K維度。代碼如下所示__global__ void matrixMulTiled(float* C, float* A, float* B, int M, int N, int K) { // 為每個(gè)線程塊聲明共享內(nèi)存 __shared__ float ds_A[TILE_WIDTH][TILE_WIDTH]; __shared__ float ds_B[TILE_WIDTH][TILE_WIDTH]; int bx blockIdx.x, by blockIdx.y; int tx threadIdx.x, ty threadIdx.y; // 計(jì)算C中當(dāng)前線程要處理的元素坐標(biāo) int Row by * TILE_WIDTH ty; int Col bx * TILE_WIDTH tx; float Cvalue 0; // 循環(huán)遍歷所有平鋪 for (int ph 0; ph ceil(K/(float)TILE_WIDTH); ph) { // 協(xié)作加載一個(gè)平鋪的數(shù)據(jù)到共享內(nèi)存 if (Row M (ph*TILE_WIDTH tx) K) ds_A[ty][tx] A[Row * K ph * TILE_WIDTH tx]; else ds_A[ty][tx] 0.0; if (Col N (ph*TILE_WIDTH ty) K) ds_B[ty][tx] B[(ph * TILE_WIDTH ty) * N Col]; else ds_B[ty][tx] 0.0; // 等待塊內(nèi)所有線程完成加載 __syncthreads(); // 使用共享內(nèi)存中的數(shù)據(jù)計(jì)算部分和 for (int i 0; i TILE_WIDTH; i) { Cvalue ds_A[ty][i] * ds_B[i][tx]; } // 等待所有線程完成計(jì)算再進(jìn)行下一輪加載避免數(shù)據(jù)競爭 __syncthreads(); } // 將結(jié)果寫回全局內(nèi)存 if (Row M Col N) C[Row * N Col] Cvalue; }這個(gè)內(nèi)核需要以二維的塊和網(wǎng)格啟動(dòng)。通過這種方式對全局內(nèi)存的訪問量從O(MNK)降到了O(MNK / TILE_WIDTH)因?yàn)槊總€(gè)數(shù)據(jù)元素從全局內(nèi)存只加載一次到共享內(nèi)存然后被重用了TILE_WIDTH次。4. 性能分析與優(yōu)化實(shí)踐超越樣例代碼完成基本功能后作業(yè)的加分項(xiàng)往往在于性能優(yōu)化和深入分析。這里有幾個(gè)可以深挖的方向。4.1 性能測量與瓶頸定位不要憑感覺說“快了”。一定要用CUDA事件Event來精確測量內(nèi)核執(zhí)行時(shí)間cudaEvent_t start, stop; cudaEventCreate(start); cudaEventCreate(stop); cudaEventRecord(start); // 啟動(dòng)你的內(nèi)核 matrixMulKernelgrid, block(...); cudaEventRecord(stop); cudaEventSynchronize(stop); float milliseconds 0; cudaEventElapsedTime(milliseconds, start, stop); printf(Kernel time: %f ms\n, milliseconds);對比不同實(shí)現(xiàn)如樸素版本、共享內(nèi)存版本的時(shí)間。同時(shí)使用nvprof舊版或nsys新版性能分析器。它們能告訴你內(nèi)核的占用率Occupancy、全局內(nèi)存讀寫效率、共享內(nèi)存使用情況等。例如如果分析器顯示“Global Memory Load Efficiency”很低說明你的全局內(nèi)存訪問模式很差沒有合并。4.2 進(jìn)階優(yōu)化技巧嘗試在共享內(nèi)存平鋪的基礎(chǔ)上還可以嘗試以下優(yōu)化并在報(bào)告中分析效果循環(huán)展開Loop Unrolling在計(jì)算部分和的內(nèi)部循環(huán)中手動(dòng)展開幾次可以減少循環(huán)開銷和增加指令級并行。CUDA編譯器也支持#pragma unroll指令。使用向量化內(nèi)存操作如果數(shù)據(jù)是float2或float4類型可以使用向量化加載/存儲指令一次傳輸更多數(shù)據(jù)提高內(nèi)存帶寬利用率。調(diào)整線程塊大小Block Size線程塊大小如16x1625632x8256會(huì)影響占用率和共享內(nèi)存庫沖突Bank Conflict。共享內(nèi)存被組織成多個(gè)庫通常是32個(gè)如果同一個(gè)時(shí)鐘周期內(nèi)線程束Warp中多個(gè)線程訪問同一個(gè)庫的不同地址就會(huì)發(fā)生庫沖突導(dǎo)致串行化訪問。通過調(diào)整數(shù)據(jù)在共享內(nèi)存中的存儲方式如使用padding或調(diào)整線程塊維度可以緩解沖突。嘗試使用只讀數(shù)據(jù)緩存Read-Only Cache對于不變的數(shù)據(jù)如矩陣乘法中的B矩陣可以使用__ldg()指令或通過const __restrict__修飾指針引導(dǎo)編譯器使用只讀數(shù)據(jù)緩存這有時(shí)比使用L1緩存更好。4.3 與標(biāo)準(zhǔn)庫的對比一個(gè)非常有說服力的分析是將你優(yōu)化的CUDA版本與高度優(yōu)化的CPU庫如Intel MKL、OpenBLAS以及CUDA自帶的庫如cuBLAS進(jìn)行性能對比。用cuBLAS的cublasSgemm函數(shù)作為一個(gè)性能基準(zhǔn)。你會(huì)發(fā)現(xiàn)即使你用了共享內(nèi)存可能仍然遠(yuǎn)不如cuBLAS因?yàn)樗€使用了更高級的技巧如雙緩沖Double Buffering、異步拷貝、張量核心Tensor Core等。在作業(yè)報(bào)告中分析這個(gè)差距的原因能體現(xiàn)你的思考深度。5. 常見錯(cuò)誤調(diào)試與作業(yè)報(bào)告撰寫心得最后分享一些調(diào)試和完成作業(yè)報(bào)告的經(jīng)驗(yàn)。5.1 那些讓人頭疼的運(yùn)行時(shí)錯(cuò)誤“an illegal instruction was encountered”這通常也是計(jì)算能力不匹配導(dǎo)致的。確保用-archsm_xx編譯選項(xiàng)指定正確的架構(gòu)例如-archsm_89對應(yīng)RTX 40系列。可以用nvcc -archsm_xx code.cu編譯?！癱udaErrorLaunchTimeout”在Windows顯示模式下如果內(nèi)核運(yùn)行時(shí)間過長通常超過2秒WDDM驅(qū)動(dòng)會(huì)認(rèn)為顯卡失去響應(yīng)從而終止內(nèi)核。這在進(jìn)行大規(guī)模測試時(shí)可能遇到。解決方法是在Linux下運(yùn)行或在Windows下使用TCC驅(qū)動(dòng)模式僅限Tesla等計(jì)算卡或者將大任務(wù)拆分成多個(gè)短時(shí)間內(nèi)核啟動(dòng)。共享內(nèi)存使用超限每個(gè)線程塊能使用的共享內(nèi)存有限如48KB。如果你聲明__shared__ float arr[1024][1024]這顯然就超了。需要根據(jù)塊大小和數(shù)據(jù)類型精確計(jì)算。5.2 調(diào)試方法printf與cuda-gdbCUDA調(diào)試不像CPU那么方便。最樸素的調(diào)試方法是使用printf。在計(jì)算能力7.0及以上的GPU上內(nèi)核中可以直接使用printf輸出會(huì)在所有線程執(zhí)行完后顯示在控制臺。對于更復(fù)雜的問題可以使用cuda-gdbLinux或Nsight VSEWindows進(jìn)行圖形化調(diào)試可以設(shè)置斷點(diǎn)、查看變量、檢查線程狀態(tài)。5.3 撰寫一份有深度的作業(yè)報(bào)告作業(yè)報(bào)告不是代碼的復(fù)述。它應(yīng)該包含設(shè)計(jì)思路清晰說明你的線程網(wǎng)格和塊是如何劃分的為什么這么劃分考慮數(shù)據(jù)規(guī)模、硬件限制。內(nèi)存優(yōu)化策略詳細(xì)解釋你是如何使用共享內(nèi)存、常量內(nèi)存的如何解決可能存在的庫沖突。性能分析提供不同版本樸素、優(yōu)化的詳細(xì)性能數(shù)據(jù)表格。用圖表展示隨著矩陣規(guī)模增大加速比的變化。分析性能瓶頸是內(nèi)存帶寬限制還是計(jì)算限制。正確性驗(yàn)證如何驗(yàn)證結(jié)果正確與CPU計(jì)算結(jié)果對比計(jì)算相對誤差。遇到的問題與解決方案把你在環(huán)境配置、編碼、調(diào)試中踩的坑和解決方法寫出來這是報(bào)告最出彩的部分。總結(jié)與展望你的實(shí)現(xiàn)還有哪些不足如果時(shí)間允許下一步可以從哪些方向優(yōu)化如使用動(dòng)態(tài)共享內(nèi)存、嘗試CUDA Graph、利用Tensor Core完成這份作業(yè)的過程痛苦和成就感是并存的。當(dāng)你第一次看到自己編寫的CUDA內(nèi)核正確運(yùn)行并帶來可觀的加速時(shí)當(dāng)你通過調(diào)整一個(gè)參數(shù)讓性能提升10%時(shí)你會(huì)對“高性能計(jì)算”這四個(gè)字有完全不同的、更深刻的理解。這不僅僅是調(diào)用一個(gè)庫而是真正在駕馭硬件這種感覺是之前純CPU編程很難帶來的。