CUDA是並行計算的平台和類C編程模型,我們能很容易的實現並行算法,就像寫C代碼一樣。只要配備的NVIDIA GPU,就可以在許多設備上運行你的並行程序,無論是台式機、筆記本抑或平板電腦。熟悉C語言可以幫助你盡快掌握CUDA。
CUDA編程允許你的程序執行在異構系統上,即CUP和GPU,二者有各自的存儲空間,並由PCI-Express 總線區分開。因此,我們應該先注意二者術語上的區分:
代碼中,一般用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程序的處理流程:
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的可選類型有:
具體含義很好懂,就不多做解釋了。
對於返回類型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內置變量:
一般會把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的調用格式為:
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的限制:
#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 number time_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 device int dev = 0; cudaSetDevice(dev);
// set up data size of vectors int nElem = 32; printf("Vector size %d\n", nElem);
// malloc host memory size_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 side initialData(h_A, nElem); initialData(h_B, nElem); memset(hostRef, 0, nBytes); memset(gpuRef, 0, nBytes);
// malloc device global memory float *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 device cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice); cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);
// invoke kernel at host side dim3 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 side cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
// add vector at host side for result checks sumArraysOnHost(h_A, h_B, hostRef, nElem);
// check device results checkResult(hostRef, gpuRef, nElem);
// free device global memory cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
// free host memory free(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.
代碼下載:CodeSamples.zip