1. 從“Hello, World!”到并行計(jì)算為什么我們需要CUDA函數(shù)庫如果你寫過CUDA程序大概率是從一個(gè)簡單的向量加法開始的。在CPU上我們寫個(gè)循環(huán)for (int i 0; i N; i) c[i] a[i] b[i];邏輯清晰但速度感人。然后你接觸到了CUDA知道了kernelgrid, block(...)這種神奇的語法把計(jì)算任務(wù)扔給成百上千個(gè)GPU線程去并行執(zhí)行性能瞬間飆升。那一刻的興奮每個(gè)搞GPU編程的人都經(jīng)歷過。但興奮過后現(xiàn)實(shí)問題接踵而至。你很快會(huì)發(fā)現(xiàn)那個(gè)簡單的向量加法kernel里藏著不少“坑”。比如你怎么確保每個(gè)線程訪問的全局內(nèi)存地址是正確且高效的當(dāng)數(shù)據(jù)量巨大一個(gè)block的線程數(shù)比如1024不夠用需要啟動(dòng)多個(gè)block組成的grid時(shí)線程的全局索引threadIdx.x blockIdx.x * blockDim.x這個(gè)公式會(huì)不會(huì)寫錯(cuò)更復(fù)雜一點(diǎn)的如果你想在block內(nèi)部讓線程之間交換數(shù)據(jù)、協(xié)同工作或者進(jìn)行歸約求和又該怎么辦難道每次都從頭開始推導(dǎo)、手寫這些底層邏輯嗎這就是CUDA運(yùn)行時(shí)API和常用函數(shù)的價(jià)值所在。它們不是CUDA編程最炫酷的部分但卻是構(gòu)建穩(wěn)定、高效GPU程序的基石。你可以把CUDA C/C語言本身__global__,__shared__等看作是給你提供了磚塊、水泥和鋼筋讓你能蓋房子。而CUDA運(yùn)行時(shí)API以cuda為前綴的函數(shù)和設(shè)備函數(shù)如threadIdx則是你手中的瓦刀、水平儀和腳手架。沒有后者你或許也能勉強(qiáng)把磚塊壘起來但房子蓋得慢、容易歪、還可能塌。這些“工具函數(shù)”封裝了GPU硬件操作的復(fù)雜性提供了內(nèi)存管理、線程組織、設(shè)備查詢、錯(cuò)誤處理等基礎(chǔ)設(shè)施讓我們能更專注于計(jì)算邏輯本身而不是陷入硬件細(xì)節(jié)的泥潭。本文將深入CUDA常用函數(shù)的第二篇章聚焦于那些在線程索引計(jì)算、內(nèi)存操作同步、以及原子操作等核心場景中高頻出現(xiàn)的函數(shù)與用法。我會(huì)結(jié)合多年在圖像處理、科學(xué)計(jì)算項(xiàng)目中踩過的坑不僅告訴你這些函數(shù)怎么用更會(huì)剖析為什么要這樣設(shè)計(jì)以及在什么場景下選擇哪個(gè)函數(shù)最合適。你會(huì)發(fā)現(xiàn)用好這些函數(shù)是從“能跑通CUDA程序”到“寫出工業(yè)級(jí)高效CUDA程序”的關(guān)鍵一步。2. 線程索引計(jì)算超越threadIdx.x的全局視野當(dāng)我們啟動(dòng)一個(gè)kernel時(shí)需要指定執(zhí)行配置grid, block。這里的block線程塊和grid網(wǎng)格是CUDA編程模型的核心抽象。一個(gè)block包含一組線程這些線程可以快速通過共享內(nèi)存通信和同步。一個(gè)grid包含多個(gè)block這些block可以獨(dú)立地在任意SM流多處理器上執(zhí)行。對(duì)于kernel內(nèi)的一個(gè)線程如何知道自己獨(dú)一無二的“身份”即全局索引以便處理對(duì)應(yīng)的數(shù)據(jù)呢這就需要我們正確計(jì)算線程索引。2.1 一維情況下的索引計(jì)算基礎(chǔ)但易錯(cuò)假設(shè)我們啟動(dòng)了一個(gè)一維的grid和一維的block。這是最常見的情況。// 假設(shè)數(shù)據(jù)總量為N int threads_per_block 256; int blocks_per_grid (N threads_per_block - 1) / threads_per_block; // 向上取整 myKernelblocks_per_grid, threads_per_block(...);在kernel內(nèi)部計(jì)算全局索引的標(biāo)準(zhǔn)公式是__global__ void myKernel(float *data, int N) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx N) { // 邊界檢查至關(guān)重要 // 對(duì)data[idx]進(jìn)行操作 } }這里的關(guān)鍵變量threadIdx.x: 當(dāng)前線程在其所屬block內(nèi)的索引0 到blockDim.x-1。blockIdx.x: 當(dāng)前block在grid中的索引0 到gridDim.x-1。blockDim.x: 每個(gè)block的線程數(shù)量即我們傳入的threads_per_block。gridDim.x:grid中block的數(shù)量即我們計(jì)算的blocks_per_grid。注意邊界檢查if (idx N)是必須的因?yàn)槲覀兊腷locks_per_grid是向上取整計(jì)算的最后一個(gè)block很可能沒有完全“填滿”線程。如果不做檢查多出來的線程會(huì)訪問到數(shù)組data之外的內(nèi)存導(dǎo)致未定義行為可能是靜默的數(shù)據(jù)損壞也可能是直接報(bào)錯(cuò)。2.2 二維與三維索引計(jì)算處理圖像和體積數(shù)據(jù)的利器對(duì)于圖像處理2D數(shù)組或體渲染3D體數(shù)據(jù)使用二維或三維的grid和block更為直觀。啟動(dòng)配置示例處理一幅 width * height 的圖像dim3 block(16, 16); // 一個(gè)block有16x16256個(gè)線程 dim3 grid((width block.x - 1) / block.x, (height block.y - 1) / block.y); processImagegrid, block(d_image, width, height);Kernel內(nèi)的索引計(jì)算__global__ void processImage(float* image, int width, int height) { int x threadIdx.x blockIdx.x * blockDim.x; int y threadIdx.y blockIdx.y * blockDim.y; if (x width y height) { int index y * width x; // 將2D坐標(biāo)轉(zhuǎn)換為1D內(nèi)存線性索引 // 對(duì) image[index] 進(jìn)行操作 } }這里threadIdx,blockIdx,blockDim,gridDim都變成了dim3類型的結(jié)構(gòu)體可以通過.x,.y,.z成員訪問各維度分量。為什么選擇16x16的block這涉及到GPU硬件的一個(gè)核心約束warp。NVIDIA GPU的基本執(zhí)行單位是warp目前通常是32個(gè)線程。為了達(dá)到最佳性能一個(gè)block的線程總數(shù)最好是32的倍數(shù)。16x16256是32的8倍是一個(gè)常見的選擇。同時(shí)block的維度如16, 16, 1也會(huì)影響共享內(nèi)存的bank沖突、全局內(nèi)存合并訪問等需要根據(jù)具體算法調(diào)整。2.3 一個(gè)實(shí)戰(zhàn)中的高級(jí)技巧使用blockIdx和threadIdx進(jìn)行數(shù)據(jù)分塊有時(shí)每個(gè)線程處理一個(gè)數(shù)據(jù)元素如上面的圖像處理可能不是最高效的特別是當(dāng)每個(gè)元素的計(jì)算量很輕時(shí)內(nèi)存訪問的開銷會(huì)成為瓶頸。這時(shí)我們可以讓每個(gè)線程處理多個(gè)數(shù)據(jù)元素。假設(shè)數(shù)據(jù)量N很大我們啟動(dòng)較少的block但讓每個(gè)線程處理一個(gè)數(shù)據(jù)塊tile。__global__ void kernelProcessTile(float *data, int N, int tile_size) { // 計(jì)算該線程負(fù)責(zé)的起始索引 int start_idx (blockIdx.x * blockDim.x threadIdx.x) * tile_size; // 每個(gè)線程循環(huán)處理tile_size個(gè)元素 for (int i 0; i tile_size; i) { int idx start_idx i; if (idx N) { // 處理data[idx] } } }這種模式增加了每個(gè)線程的計(jì)算強(qiáng)度Arithmetic Intensity有助于隱藏內(nèi)存訪問延遲提升整體吞吐量。在優(yōu)化卷積、矩陣乘法等核心時(shí)這種“分塊”思想是基礎(chǔ)。3. 線程協(xié)作與同步讓block內(nèi)的線程“步調(diào)一致”CUDA的線程層次結(jié)構(gòu)中block內(nèi)的線程可以通過共享內(nèi)存Shared Memory進(jìn)行高速通信并通過同步函數(shù)來協(xié)調(diào)執(zhí)行順序。這是實(shí)現(xiàn)很多高效并行算法如歸約、掃描、卷積的關(guān)鍵。3.1__syncthreads()block內(nèi)的柵欄__syncthreads()是一個(gè)CUDA內(nèi)置函數(shù)intrinsic function用于同步同一個(gè)block內(nèi)的所有線程。當(dāng)某個(gè)線程調(diào)用它時(shí)它會(huì)等待直到該block內(nèi)的所有線程都執(zhí)行到這個(gè)調(diào)用點(diǎn)然后所有線程才會(huì)繼續(xù)執(zhí)行之后的代碼。一個(gè)典型應(yīng)用場景歸約求和Reduction假設(shè)我們要計(jì)算一個(gè)block內(nèi)所有線程的某個(gè)局部變量的總和。__global__ void sumReduction(float *input, float *output) { // 聲明共享內(nèi)存大小等于一個(gè)block的線程數(shù) extern __shared__ float s_data[]; unsigned int tid threadIdx.x; unsigned int i blockIdx.x * blockDim.x threadIdx.x; // 每個(gè)線程將全局內(nèi)存數(shù)據(jù)加載到共享內(nèi)存 s_data[tid] input[i]; // 等待所有線程完成加載這是必須的。 __syncthreads(); // 在共享內(nèi)存上進(jìn)行歸約求和 for (unsigned int s blockDim.x / 2; s 0; s 1) { if (tid s) { s_data[tid] s_data[tid s]; } // 等待上一輪加法完成確保數(shù)據(jù)就緒再進(jìn)行下一輪 __syncthreads(); } // 現(xiàn)在總和在s_data[0]中 if (tid 0) { output[blockIdx.x] s_data[0]; } }為什么這里需要多次__syncthreads()第一次同步加載后確保所有線程都已將自己負(fù)責(zé)的數(shù)據(jù)寫入共享內(nèi)存s_data的對(duì)應(yīng)位置。如果沒有這個(gè)同步某個(gè)線程可能還在加載數(shù)據(jù)而其他線程已經(jīng)開始讀取s_data進(jìn)行求和導(dǎo)致讀到的是未初始化的舊值。循環(huán)內(nèi)的同步每輪歸約后每一輪歸約操作s_data[tid] s_data[tid s]都修改了共享內(nèi)存。必須等待這一輪所有參與計(jì)算的線程都完成了寫入操作下一輪讀取s_data的線程才能獲得正確的結(jié)果。否則會(huì)發(fā)生數(shù)據(jù)競爭Data Race。警告__syncthreads()的使用必須非常小心它要求同一個(gè)block內(nèi)的所有線程都必須執(zhí)行到這個(gè)調(diào)用點(diǎn)否則程序會(huì)死鎖。這意味著在條件分支中如果只有部分線程執(zhí)行了__syncthreads()而其他線程因?yàn)闂l件不滿足跳過了它GPU就會(huì)一直等待那些永遠(yuǎn)不會(huì)到達(dá)的線程導(dǎo)致kernel掛起。因此確保__syncthreads()在代碼路徑中是無條件執(zhí)行的或者所有線程都經(jīng)過完全相同的條件分支。3.2 內(nèi)存屏障__threadfence()系列函數(shù)__syncthreads()只保證block內(nèi)部的線程看到一致的共享內(nèi)存和全局內(nèi)存狀態(tài)。但在某些涉及多個(gè)block協(xié)作或者需要確保對(duì)全局內(nèi)存的寫入對(duì)其他線程尤其是其他block的線程或主機(jī)線程可見時(shí)就需要更強(qiáng)的內(nèi)存排序保證。這時(shí)就需要內(nèi)存屏障Memory Fence。CUDA提供了不同粒度的內(nèi)存屏障函數(shù)__threadfence(): 確保調(diào)用線程在屏障之前對(duì)所有內(nèi)存全局、共享、本地的寫入在該屏障之后對(duì)同一設(shè)備上的所有線程可見。__threadfence_block(): 確保調(diào)用線程在屏障之前對(duì)所有內(nèi)存的寫入在該屏障之后對(duì)同一block內(nèi)的所有線程可見。它比__syncthreads()弱不要求線程同步只要求內(nèi)存操作順序。__threadfence_system(): 功能最強(qiáng)確保寫入對(duì)所有線程包括主機(jī)線程可見。這通常用于支持全局原子操作或與主機(jī)端進(jìn)行信號(hào)量等同步。一個(gè)使用場景示例生產(chǎn)者-消費(fèi)者模型假設(shè)有多個(gè)block向全局內(nèi)存的一個(gè)隊(duì)列寫入數(shù)據(jù)另一個(gè)block或主機(jī)從隊(duì)列中讀取。生產(chǎn)者block在寫入數(shù)據(jù)后需要確保數(shù)據(jù)真正寫入了全局內(nèi)存并且寫入操作對(duì)其他線程可見然后才能更新隊(duì)列的“尾指針”。這個(gè)更新指針的操作就需要用__threadfence()或原子操作來保護(hù)以防止消費(fèi)者讀到未完全寫入的數(shù)據(jù)。// 生產(chǎn)者線程簡化示意 __global__ void producer(int *queue, int *tail_index, int data) { // ... 計(jì)算寫入位置 my_index ... queue[my_index] data; // 寫入數(shù)據(jù) // 確保數(shù)據(jù)寫入對(duì)所有人可見 __threadfence(); // 然后才能安全地更新尾指針這通常需要一個(gè)原子操作 // atomicAdd(tail_index, 1); }在實(shí)際項(xiàng)目中__threadfence()的使用相對(duì)高階通常與原子操作、鎖、信號(hào)量等同步原語配合使用。對(duì)于大部分初學(xué)者和常見算法__syncthreads()已經(jīng)足夠。4. 原子操作在多線程中安全地“爭搶”資源當(dāng)多個(gè)線程需要讀寫同一個(gè)內(nèi)存位置時(shí)就會(huì)發(fā)生數(shù)據(jù)競爭。例如多個(gè)線程同時(shí)對(duì)一個(gè)全局計(jì)數(shù)器進(jìn)行count操作。在CPU上我們可以用互斥鎖mutex。在GPU上鎖的實(shí)現(xiàn)成本很高容易導(dǎo)致線程束內(nèi)線程分化嚴(yán)重降低性能。CUDA提供了更輕量級(jí)的解決方案原子操作Atomic Operations。原子操作保證了對(duì)某個(gè)內(nèi)存地址的“讀-修改-寫”操作是不可分割的。在執(zhí)行過程中不會(huì)有其他線程能訪問到該內(nèi)存地址的中間狀態(tài)。4.1 常用原子函數(shù)一覽CUDA C 提供了豐富的原子函數(shù)它們定義在cuda/atomic頭文件中CUDA 11及以上推薦使用C標(biāo)準(zhǔn)風(fēng)格的原子操作也有傳統(tǒng)的內(nèi)置函數(shù)。這里以傳統(tǒng)函數(shù)為例因?yàn)樗鼈兏Ｒ娪诂F(xiàn)有代碼。atomicAdd(int* address, int val): 原子加法。*address val。atomicSub(int* address, int val): 原子減法。atomicExch(int* address, int val): 原子交換。將*address的值設(shè)置為val并返回舊值。atomicMin(int* address, int val),atomicMax: 原子最小/最大值。atomicInc,atomicDec: 原子遞增/遞減帶環(huán)繞。atomicCAS(int* address, int compare, int val):比較并交換Compare And Swap。這是最強(qiáng)大、最基礎(chǔ)的原子操作。如果*address compare則將*address設(shè)置為val否則不修改。無論是否修改都返回*address的舊值。它可以用來構(gòu)建任何其他原子操作如鎖。支持的數(shù)據(jù)類型上述函數(shù)通常有對(duì)應(yīng)unsigned int,unsigned long long int,float(僅atomicAdd等部分操作),double(計(jì)算能力6.0) 的版本。4.2 原子操作的性能代價(jià)與使用策略原子操作雖然方便但代價(jià)高昂。因?yàn)楫?dāng)多個(gè)線程同時(shí)原子訪問同一個(gè)內(nèi)存地址時(shí)硬件會(huì)將這些操作串行化導(dǎo)致嚴(yán)重的性能下降。這被稱為競爭Contention。使用原則能不用則不用首先考慮算法是否可以重構(gòu)避免對(duì)共享資源的競爭。例如使用歸約讓每個(gè)block先計(jì)算局部和再由一個(gè)線程將局部和加到全局變量上這樣全局原子操作的次數(shù)就從數(shù)據(jù)量N減少到了block的數(shù)量競爭大大降低。分散競爭如果必須使用原子操作盡量讓線程原子訪問不同的內(nèi)存地址。例如在直方圖統(tǒng)計(jì)中如果直接讓所有線程原子增加一個(gè)全局直方圖數(shù)組競爭會(huì)非常激烈。更好的方法是使用私有化Privatization每個(gè)block先在共享內(nèi)存中構(gòu)建一個(gè)局部直方圖這個(gè)過程可能也需要原子操作但共享內(nèi)存的原子操作速度遠(yuǎn)快于全局內(nèi)存計(jì)算完成后再由一個(gè)線程將局部直方圖累加到全局內(nèi)存中。這樣全局原子操作的次數(shù)從像素?cái)?shù)減少到了直方圖的桶數(shù)bin數(shù)。使用更快的原子操作空間CUDA提供了在共享內(nèi)存上進(jìn)行原子操作的函數(shù)如atomicAdd在共享內(nèi)存上同樣可用。共享內(nèi)存的延遲和帶寬遠(yuǎn)優(yōu)于全局內(nèi)存因此競爭開銷小得多。4.3 實(shí)戰(zhàn)案例使用原子操作實(shí)現(xiàn)簡單的唯一ID分配器假設(shè)我們有一個(gè)任務(wù)隊(duì)列需要為每個(gè)新任務(wù)分配一個(gè)全局唯一的ID。__device__ int global_next_id 0; // 存儲(chǔ)在全局內(nèi)存中的設(shè)備變量 __global__ void generateTaskId(int *task_ids, int num_tasks) { int idx threadIdx.x blockIdx.x * blockDim.x; if (idx num_tasks) { // 每個(gè)線程原子地獲取并增加全局ID計(jì)數(shù)器 int my_id atomicAdd(global_next_id, 1); task_ids[idx] my_id; } }這個(gè)kernel啟動(dòng)后task_ids數(shù)組將被填入0, 1, 2, ... 這樣唯一的ID。atomicAdd確保了即使成千上萬個(gè)線程同時(shí)執(zhí)行每個(gè)線程得到的my_id也是不同的。潛在問題如果num_tasks非常大所有線程都競爭同一個(gè)全局變量global_next_id性能會(huì)很差。在實(shí)際生產(chǎn)中更優(yōu)的方案可能是每個(gè)block預(yù)先原子獲取一個(gè)ID范圍比如atomicAdd(global_next_id, blockDim.x)然后在block內(nèi)部線性分配。這減少了原子操作次數(shù)。或者直接使用blockIdx.x * blockDim.x threadIdx.x這種計(jì)算出來的索引作為唯一ID如果ID不需要全局嚴(yán)格連續(xù)的話。5. 內(nèi)存操作函數(shù)高效地在GPU與GPU、GPU與CPU間搬運(yùn)數(shù)據(jù)除了內(nèi)核函數(shù)CUDA運(yùn)行時(shí)API中另一大類常用函數(shù)就是內(nèi)存操作函數(shù)。高效、正確地管理內(nèi)存是GPU編程性能的關(guān)鍵。5.1cudaMemcpy家族數(shù)據(jù)搬運(yùn)的主力cudaMemcpy用于在主機(jī)Host內(nèi)存和設(shè)備Device內(nèi)存之間或者設(shè)備內(nèi)存內(nèi)部復(fù)制數(shù)據(jù)。它的基本原型是cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);其中cudaMemcpyKind指定了復(fù)制方向cudaMemcpyHostToHost: 主機(jī)到主機(jī)一般不用。cudaMemcpyHostToDevice: 主機(jī)到設(shè)備。將數(shù)據(jù)從CPU內(nèi)存拷貝到GPU顯存。cudaMemcpyDeviceToHost: 設(shè)備到主機(jī)。將計(jì)算結(jié)果從GPU顯存拷貝回CPU內(nèi)存。cudaMemcpyDeviceToDevice: 設(shè)備內(nèi)部拷貝。在兩個(gè)GPU顯存區(qū)域之間拷貝。使用示例float *h_data new float[N]; // 主機(jī)內(nèi)存 float *d_data; cudaMalloc(d_data, N * sizeof(float)); // 設(shè)備內(nèi)存 // 初始化主機(jī)數(shù)據(jù) for (int i 0; i N; i) h_data[i] i; // 主機(jī) - 設(shè)備 cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice); // ... 在GPU上執(zhí)行kernel處理d_data ... // 設(shè)備 - 主機(jī) cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost); // 清理 cudaFree(d_data); delete[] h_data;5.2 異步內(nèi)存拷貝cudaMemcpyAsync與流管理標(biāo)準(zhǔn)的cudaMemcpy是同步的調(diào)用會(huì)阻塞主機(jī)線程直到拷貝完成。這對(duì)于小數(shù)據(jù)量沒問題但對(duì)于大數(shù)據(jù)量或需要重疊計(jì)算與傳輸?shù)膱鼍熬托枰惒娇截恈udaMemcpyAsync。異步拷貝需要與CUDA流Stream配合使用。流是一系列順序執(zhí)行的命令如內(nèi)存拷貝、內(nèi)核啟動(dòng)的隊(duì)列。不同流中的命令可以并發(fā)執(zhí)行如果硬件支持。cudaStream_t stream; cudaStreamCreate(stream); // 創(chuàng)建一個(gè)流 float *h_pinned_data; cudaMallocHost(h_pinned_data, N * sizeof(float)); // 分配鎖頁主機(jī)內(nèi)存Page-Locked Host Memory // 異步拷貝主機(jī)-設(shè)備該操作被放入流中立即返回 cudaMemcpyAsync(d_data, h_pinned_data, N * sizeof(float), cudaMemcpyHostToDevice, stream); // 主機(jī)線程可以繼續(xù)做其他事情不必等待拷貝完成 // 在同一個(gè)流中啟動(dòng)kernel它會(huì)等待之前的拷貝完成 myKernelgrid, block, 0, stream(d_data, N); // 異步拷貝結(jié)果回主機(jī) cudaMemcpyAsync(h_pinned_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost, stream); // 等待流中的所有操作完成 cudaStreamSynchronize(stream); // 清理 cudaStreamDestroy(stream); cudaFreeHost(h_pinned_data);關(guān)鍵點(diǎn)鎖頁內(nèi)存cudaMemcpyAsync的源或目標(biāo)地址如果是主機(jī)內(nèi)存則必須是鎖頁內(nèi)存通過cudaMallocHost分配或使用cudaHostAlloc。普通malloc分配的內(nèi)存是可分頁的DMA引擎無法安全地異步訪問。并發(fā)性通過創(chuàng)建多個(gè)流并將內(nèi)存拷貝和內(nèi)核執(zhí)行任務(wù)分配到不同流中可以實(shí)現(xiàn)計(jì)算與數(shù)據(jù)傳輸?shù)闹丿B這是提升GPU程序整體吞吐量的重要手段。cudaStreamSynchronize(stream): 阻塞主機(jī)線程直到指定流中的所有任務(wù)完成。5.3 統(tǒng)一內(nèi)存Unified Memory與cudaMallocManaged從CUDA 6.0開始引入了統(tǒng)一內(nèi)存Unified Memory, UM模型。它提供了一個(gè)統(tǒng)一的內(nèi)存地址空間從CPU和GPU都可以訪問。內(nèi)存的物理位置由CUDA驅(qū)動(dòng)和運(yùn)行時(shí)自動(dòng)管理在需要時(shí)進(jìn)行遷移。這極大地簡化了編程。// 使用統(tǒng)一內(nèi)存分配 float *um_data; cudaMallocManaged(um_data, N * sizeof(float)); // CPU可以直接訪問和初始化 for (int i 0; i N; i) um_data[i] i; // 啟動(dòng)kernelGPU訪問數(shù)據(jù)。運(yùn)行時(shí)會(huì)在首次訪問時(shí)自動(dòng)將數(shù)據(jù)遷移到GPU。 myKernelgrid, block(um_data, N); cudaDeviceSynchronize(); // 等待kernel完成 // CPU可以立即讀取結(jié)果數(shù)據(jù)會(huì)被自動(dòng)遷回。 printf(%f\n, um_data[0]); cudaFree(um_data);優(yōu)點(diǎn)編程簡單無需手動(dòng)cudaMemcpy。對(duì)于復(fù)雜的數(shù)據(jù)結(jié)構(gòu)如鏈表、樹在CPU和GPU間共享特別有用。缺點(diǎn)存在一定的性能開銷頁面遷移、缺頁中斷。對(duì)于規(guī)律性的大規(guī)模數(shù)據(jù)搬運(yùn)手動(dòng)管理內(nèi)存拷貝和流通常能獲得更優(yōu)的性能。統(tǒng)一內(nèi)存更適合于原型開發(fā)、簡化代碼或者訪問模式不規(guī)律、數(shù)據(jù)依賴復(fù)雜的場景。6. 錯(cuò)誤處理給你的CUDA代碼加上“安全帶”CUDA API函數(shù)和內(nèi)核啟動(dòng)幾乎都會(huì)返回一個(gè)類型為cudaError_t的錯(cuò)誤碼。忽略錯(cuò)誤檢查是CUDA新手最常見的錯(cuò)誤之一這會(huì)導(dǎo)致程序在出現(xiàn)問題時(shí)行為詭異難以調(diào)試。6.1 檢查運(yùn)行時(shí)API錯(cuò)誤每個(gè)CUDA運(yùn)行時(shí)API調(diào)用后都應(yīng)該檢查錯(cuò)誤?？梢詫懸粋€(gè)簡單的宏來包裝#define CHECK(call) \ { \ const cudaError_t error call; \ if (error ! cudaSuccess) { \ printf(Error: %s:%d, , __FILE__, __LINE__); \ printf(code:%d, reason: %s\n, error, cudaGetErrorString(error)); \ exit(1); \ } \ } // 使用方式 CHECK(cudaMalloc(d_data, N * sizeof(float))); CHECK(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice));cudaGetErrorString(error)可以將錯(cuò)誤碼轉(zhuǎn)換為可讀的描述信息。6.2 檢查內(nèi)核啟動(dòng)錯(cuò)誤內(nèi)核啟動(dòng)是異步的調(diào)用kernel...會(huì)立即返回。如果配置錯(cuò)誤如線程數(shù)超限、共享內(nèi)存分配過多這個(gè)錯(cuò)誤并不會(huì)立即被捕獲。為了捕獲內(nèi)核啟動(dòng)錯(cuò)誤需要同步設(shè)備或使用cudaGetLastError。myKernelgrid, block(...); // 方法1同步設(shè)備這會(huì)等待kernel完成并返回可能的錯(cuò)誤 cudaError_t err cudaDeviceSynchronize(); if (err ! cudaSuccess) { printf(Kernel launch failed: %s\n, cudaGetErrorString(err)); } // 方法2獲取最后一個(gè)錯(cuò)誤不等待kernel完成 cudaError_t err cudaGetLastError(); if (err ! cudaSuccess) { printf(Kernel launch failed (async): %s\n, cudaGetErrorString(err)); }通常在調(diào)試階段可以在每個(gè)內(nèi)核啟動(dòng)后使用cudaDeviceSynchronize()來確保錯(cuò)誤被及時(shí)捕獲。在生產(chǎn)代碼中可能只在關(guān)鍵位置同步。6.3 更精細(xì)的錯(cuò)誤檢查cudaPeekAtLastError與cudaGetLastErrorcudaGetLastError(): 獲取最后一個(gè)錯(cuò)誤并重置錯(cuò)誤狀態(tài)為cudaSuccess。cudaPeekAtLastError(): 獲取最后一個(gè)錯(cuò)誤但不重置錯(cuò)誤狀態(tài)。這意味著如果你連續(xù)調(diào)用兩次cudaGetLastError()第二次很可能返回cudaSuccess因?yàn)榈谝淮握{(diào)用已經(jīng)清除了錯(cuò)誤。而cudaPeekAtLastError()允許你檢查錯(cuò)誤而不清除它這在某些復(fù)雜的錯(cuò)誤處理邏輯中可能有用。對(duì)于大多數(shù)情況使用cudaGetLastError()或直接檢查API返回值即可。養(yǎng)成檢查每一個(gè)CUDA調(diào)用返回值的習(xí)慣雖然會(huì)讓代碼看起來冗長一些但在程序出錯(cuò)時(shí)它能為你節(jié)省數(shù)小時(shí)甚至數(shù)天的調(diào)試時(shí)間。這是寫出健壯CUDA程序的基石。