CUDA學習（二）

程式設計模型：

核函數Kernel：

CUDA C通過允許程式員定義C函數（稱為核心）來擴充C，C函數在被調用時被N個不同的CUDA線程并行執行N次，而不是像正常C函數那樣隻執行一次。

一個核函數使用_global_聲明定義，并使用<<<…>>>配置該核函數的CUDA線程數。執行核心的每個線程都有一個唯一的線程ID，可通過内置的threadIdx變量在核心中通路。

下面是一個例子（部分代碼）：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
...
}           

這裡，執行VecAdd（）的N個線程中的每一個執行一對成對的加法。

線程結構Thread Hierarchy

設threadIdx是一個3維的向量，可以用一維，二維或者三維的線程索引來辨別線程，形成一維，二維和三維的塊線程，稱為線程塊。

線程的索引及其線程ID以直接的方式互相關聯：對于一維的線程塊，他們是相同的；對于二維的線程塊大小$(D_{x},D_{y})$,線程$(x,y)$的索引為$(x+y*D_{x})$；對于三維線程塊$(D_{x},D_{y},D_{z})$，線程$(x,y,z)$的線程索引為$(x+y*D_{x}+z*D_{y}*D_{z})$

下面是一個矩陣A乘B，N*N的例子：

:
// 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);
...
}           

每塊的線程數是有限制的，因為一個塊的所有線程都應該駐留在同一個處理器核心上，并且必須共享該核心的有限記憶體資源。 在目前的GPU上，一個線程塊最多可以包含1024個線程

但是，核心可以由多個相同形狀的線程塊執行，是以線程的總數等于每塊的線程數乘以塊數

下圖顯示Blocks被組織成1維，2維，3維的線程塊形式.。

每個塊的線程數和<<< ... >>>文法中指定的每個網格的塊數可以是int類型或dim3類型。

網格中的每個塊都可以通過核心中通過内置的blockIdx變量在核心中通路的一維，二維或三維索引來辨別。 線程塊的次元可以通過内置的blockDim變量在核心中通路。

下面是關于block形式下的矩陣相加：

.
// 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個線程）是常見的選擇。網格（grid）是由blocks組成的，每個矩陣的元素是一個線程。這個例子假設每個次元中每個網格的線程數量可以被該次元中每個塊的線程數量平均分割，盡管不一定是這種情況。

線程塊需要獨立執行：必須能夠以任意順序，并行或串行執行它們。

塊内的線程可以通過共享記憶體共享資料并同步執行來協調記憶體通路；

更确切地說，可以通過調用————syncthreads（）内部函數來同步同一個核心中的線程;_syncthreads()是一個栅欄，所有的線程都必須到達這個栅欄時才允許線程繼續下去。 除__syncthreads（）之外，協作組API還提供了一組豐富的線程同步基元。

為了有效的協作，共享記憶體應該是每個處理器核心附近的低延遲記憶體（很像L1緩存），__syncthreads（）預計會是輕量級的。