CUDA C編程入門-編程模型

標籤：style   blog   http   color   使用   strong   io   資料   

　　這章節介紹CUDA編程模型的主要的概念。

2.1.kernels（核函數）

　　CUDA C擴充了C語言，允許程式員定義C函數，稱為kernels（核函數）。並行地在N個CUDA線程中執行N次。

　　使用__global__說明符聲明一個核函數，調用使用<<<...>>>，並且指定執行的CUDA線程數目。執行的每個線程都有一個獨一的ID，在核函數中可以通過變數threadIdx擷取。

　　例子，兩個向量的加，A加B，並把結果存入C，A、B和C的長度為N。

__global__ void addKernel(int *c, const int *a, const int *b){    int i = threadIdx.x;    c[i] = a[i] + b[i];}int main(){  ...  // Launch a kernel on the GPU with one thread for each element.    addKernel<<<1, N>>>(c, a, b);  ...      }

 

　　其中，每一個線程都會在數組中的每個元素上執行addKernel這個核函數。

2.2.線程層次

　　threadIdx是一個3元組，因此線程可以被一維、二維和三維的threadIdx標識，形成一維、二維和三維的線程塊。

　　線程的索引和ID之間的關係：對於一維的線程塊，索引和ID是相同的；對於大小為(Dx,Dy)的二維的線程塊，索引為(x,y)，而ID為x+y*Dx；對於大小為(Dx,Dy,Dz)的三維線程塊，索引為(x,y,z)，ID為(x+y*Dx+z*Dx*Dy)。

　　例子，二維矩陣加，N*N大小的A和B相加，結果存入C。

　　

// Kernel definition__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){    int i = threadIdx.x;    int j = threadIdx.y;    C[i][j] = A[i][j] + B[i][j];}int main(){    ...    // Kernel invocation with one block of N * N * 1 threads    int numBlocks = 1;    dim3 threadsPerBlock(N, N);    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);    ...}

　　其中，因為一個block的線程一般會在同一個處理核心中，並且共用有限的記憶體，所以一個塊的線程的個數是有限制的。在當前的GPU，塊中的線程數目最大為1024。

　　然而，一個核函數可以在多個相同大小的線程塊中執行，所以總的線程數等於線程塊的個數乘以每個線程塊中線程的個數。

　　線程塊被組織進一維、二維或者三維的grid中，6所示。

圖6 擁有多個線程塊的grid

　　調用核函數的時候，可以通過<<<...>>>指定每個線程塊中線程的個數以及每個grid中線程塊的個數，<<<...>>>的類型可以為int何dim3。核函數中可以通過內建的變數blockIdx得到grid中的每個線程塊的索引。同時可以通過blockDim獲得每個線程塊的維數。

　　例子，擴充先前的矩陣加的例子為多個線程塊的。

// Kernel definition__global__ void MatAdd(float A[N][N], float B[N][N],float C[N][N]){    int i = blockIdx.x * blockDim.x + threadIdx.x;    int j = blockIdx.y * blockDim.y + threadIdx.y;    if (i < N && j < N)    C[i][j] = A[i][j] + B[i][j];}int main(){    ...    // Kernel invocation    dim3 threadsPerBlock(16, 16);    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);    ...}

　　其中，線程塊的大小為16*16，總共有256個線程。在同一個線程塊中線程可以通過shared memory（共用記憶體）共用資料。可以使用__syncthreads()函數同步線程對共用記憶體的資料訪問。

2.3.記憶體層次

　　執行時，核函數可以擷取多個記憶體空間中的資料，7所示。每個線程有自己的局部記憶體。每個線程塊擁有共用記憶體空間，塊中的每個線程都可以訪問。還有執行核函數的每個線程都可以訪問的全域記憶體空間。

　　另外，還有兩個額外的唯讀記憶體空間：常量和紋理記憶體空間。全域、常量和紋理記憶體空間為不同的記憶體用法做了最佳的最佳化。紋理模型提供多種的定址方式。

 

圖7 記憶體層次

2.4.異構編程

　　8所示，CUDA編程模型假設CUDA線程執行在一個與主機的C程式分離的裝置上。核函數在GPU上執行，而其他的在CPU上執行。CUDA編程模型同時假設主機和裝置獨立操作它們在DRAM中的記憶體空間。因此，程式調用CUDA運行時（在編程介面這一章描述）管理全域、常量和紋理記憶體空間對核函數的訪問性，運行時包括記憶體配置和回收、主機和裝置之間的記憶體資料的拷貝等。

在主機上執行串列的代碼，而GPU執行並行的核函數

圖8 異構編程

2.5.計算能力

　　裝置的計算能力定義為一個主要版本號和一個次的修訂編號。

　　相同的主要版本號的GPU擁有相同的核心架構。主版號為5的是Maxwell架構，3為Kepler架構，2為Fermi架構，1為Tesla架構。

　　次修訂編號相當於核的不斷改進和新的特性。

　　支援CUDA的GPU這一章節列出支援CUDA的所有GPU的計算能力。計算能力這一章節給出每個計算能力的詳細技術規格。