CUDA簡介

CUDA是并行計算的平臺和類C編程模型，我們能很容易的實現并行算法，就像寫C代碼一樣。只要配備的NVIDIA GPU，就可以在許多設備上運行你的并行程序，無論是臺式機、筆記本抑或平板電腦。熟悉C語言可以幫助你盡快掌握CUDA。

CUDA編程

CUDA編程允許你的程序執行在異構系統上，即CUP和GPU，二者有各自的存儲空間，并由PCI-Express 總線區分開。因此，我們應該先注意二者術語上的區分：

• Host：CPU and itsmemory (host memory)
• Device: GPU and its memory (device memory)

代碼中，一般用h_前綴表示host memory，d_表示device memory。

kernel是CUDA編程中的關鍵，他是跑在GPU的代碼，用標示符__global__注明。

host可以獨立于host進行大部分操作。當一個kernel啟動后，控制權會立刻返還給CPU來執行其他額外的任務。所以，CUDA編程是異步的。一個典型的CUDA程序包含由并行代碼補足的串行代碼，串行代碼由host執行，并行代碼在device中執行。host端代碼是標準C，device是CUDA C代碼。我們可以把所有代碼放到一個單獨的源文件，也可以使用多個文件或庫。NVIDIA C編譯器（nvcc）可以編譯host和device生成可執行程序。

這里再次說明下CUDA程序的處理流程：

1. 從CPU拷貝數據到GPU。
2. 調用kernel來操作存儲在GPU的數據。
3. 將操作結果從GPU拷貝至CPU。

Memory操作

cuda程序將系統區分成host和device，二者有各自的memory。kernel可以操作device memory，為了能很好的控制device端內存，CUDA提供了幾個內存操作函數：

?

為了保證和易于學習，CUDA C 的風格跟C很接近，比如：

cudaError_t cudaMalloc ( void** devPtr, size_t size )

我們主要看看cudaMencpy，其函數原型為：

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )

其中cudaMemcpykind的可選類型有：

1. cudaMemcpyHostToHost
2. cudaMemcpyHossToDevice
3. cudaMemcpyDeviceToHost
4. cudaMemcpuDeviceToDevice

具體含義很好懂，就不多做解釋了。

對于返回類型cudaError_t，如果正確調用，則返回cudaSuccess，否則返回cudaErrorMemoryAllocation。可以使用char* cudaGetErrorString(cudaError_t error)將其轉化為易于理解的格式。

組織線程

掌握如何組織線程是CUDA編程的重要部分。CUDA線程分成Grid和Block兩個層次。

?

由一個單獨的kernel啟動的所有線程組成一個grid，grid中所有線程共享global memory。一個grid由許多block組成，block由許多線程組成，grid和block都可以是一維二維或者三維，上圖是一個二維grid和二維block。

這里介紹幾個CUDA內置變量：

• blockIdx：block的索引，blockIdx.x表示block的x坐標。
• threadIdx：線程索引，同理blockIdx。
• blockDim：block維度，上圖中blockDim.x=5.
• gridDim：grid維度，同理blockDim。

一般會把grid組織成2D，block為3D。grid和block都使用dim3作為聲明，例如：

dim3 block(3);
// 后續博文會解釋為何這樣寫grid
dim3 grid((nElem+block.x-1)/block.x);

需要注意的是，dim3僅為host端可見，其對應的device端類型為uint3。

啟動CUDA kernel

CUDA kernel的調用格式為：

kernel_name<<<grid, block>>>(argument list);

其中grid和block即為上文中介紹的類型為dim3的變量。通過這兩個變量可以配置一個kernel的線程總和，以及線程的組織形式。例如：

kernel_name<<<4, 8>>>(argumentt list);

該行代碼表明有grid為一維，有4個block，block為一維，每個block有8個線程，故此共有4*8=32個線程。

注意，不同于c函數的調用，所有CUDA kernel的啟動都是異步的，當CUDA kernel被調用時，控制權會立即返回給CPU。

函數類型標示符

__device__ 和__host__可以組合使用。

kernel的限制：

• 僅能獲取device memory 。
• 必須返回void類型。
• 不支持可變數目參數。
• 不支持靜態變量。
• 不支持函數指針。
• 異步。

代碼分析

#include <cuda_runtime.h>
#include <stdio.h>
#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); \} \
}
void checkResult(float *hostRef, float *gpuRef, const int N) {double epsilon = 1.0E-8;bool match = 1;for (int i=0; i<N; i++) {if (abs(hostRef[i] - gpuRef[i]) > epsilon) {match = 0;printf("Arrays do not match!\n");printf("host %5.2f gpu %5.2f at current %d\n",hostRef[i],gpuRef[i],i);break;}}if (match) printf("Arrays match.\n\n");
}
void initialData(float *ip,int size) {// generate different seed for random numbertime_t t;srand((unsigned) time(&t));for (int i=0; i<size; i++) {ip[i] = (float)( rand() & 0xFF )/10.0f;}
}
void sumArraysOnHost(float *A, float *B, float *C, const int N) {for (int idx=0; idx<N; idx++)C[idx] = A[idx] + B[idx];
}__global__ void sumArraysOnGPU(float *A, float *B, float *C) {int i = threadIdx.x;C[i] = A[i] + B[i];
}
int main(int argc, char **argv) {printf("%s Starting...\n", argv[0]);// set up deviceint dev = 0;cudaSetDevice(dev);// set up data size of vectorsint nElem = 32;printf("Vector size %d\n", nElem);// malloc host memorysize_t nBytes = nElem * sizeof(float);float *h_A, *h_B, *hostRef, *gpuRef;h_A = (float *)malloc(nBytes);h_B = (float *)malloc(nBytes);hostRef = (float *)malloc(nBytes);gpuRef = (float *)malloc(nBytes);// initialize data at host sideinitialData(h_A, nElem);initialData(h_B, nElem);memset(hostRef, 0, nBytes);memset(gpuRef, 0, nBytes);// malloc device global memoryfloat *d_A, *d_B, *d_C;cudaMalloc((float**)&d_A, nBytes);cudaMalloc((float**)&d_B, nBytes);cudaMalloc((float**)&d_C, nBytes);// transfer data from host to devicecudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);// invoke kernel at host sidedim3 block (nElem);dim3 grid (nElem/block.x);sumArraysOnGPU<<< grid, block >>>(d_A, d_B, d_C);printf("Execution configuration <<<%d, %d>>>\n",grid.x,block.x);// copy kernel result back to host sidecudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);// add vector at host side for result checkssumArraysOnHost(h_A, h_B, hostRef, nElem);// check device resultscheckResult(hostRef, gpuRef, nElem);// free device global memorycudaFree(d_A);cudaFree(d_B);cudaFree(d_C);// free host memoryfree(h_A);free(h_B);free(hostRef);free(gpuRef);return(0);
}

編譯指令：$nvcc sum.cu -o sum

運行： $./sum

輸出：

./sum Starting...
Vector size 32
Execution configuration <<<1, 32>>>

Arrays match.

總結

以上是生活随笔為你收集整理的CUDA学习2的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。