CUDA学习2
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程序的處理流程:
- 從CPU拷貝數據到GPU。
- 調用kernel來操作存儲在GPU的數據。
- 將操作結果從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的可選類型有:
- cudaMemcpyHostToHost
- cudaMemcpyHossToDevice
- cudaMemcpyDeviceToHost
- 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.
總結
- 上一篇: Ray Tracing,Ray Cast
- 下一篇: CUDA学习1