介紹CUDA編程模型及CUDA線程體系

可伸縮的編程模型

CUDA 編程模型主要有三個關(guān)鍵抽象：層級的線程組，共享內(nèi)存和柵同步（barrier synchronization）。

這些抽象提供了細(xì)粒度的數(shù)據(jù)并行和線程并行，可以以嵌套在粗粒的數(shù)據(jù)并行和任務(wù)并行中。它們鼓勵將問題分解為子問題。每個子問題可以獨(dú)立的在block threads中并行解決。同時每個子問題分成更細(xì)的部分，可以由塊中的所有線程并行地合作解決。

這種分解通過允許線程在解決每個子問題時進(jìn)行協(xié)作來保留語言的表達(dá)性，同時支持自動可伸縮性。實(shí)際上，每個線程塊都可以在GPU中任何可用的多處理器上調(diào)度，以任何順序、并發(fā)或順序，因此編譯的CUDA程序可以在任意數(shù)量的多處理器上執(zhí)行，如圖所示，而且只有運(yùn)行時系統(tǒng)需要知道物理多處理器的數(shù)量。

圖1 Automatic Scalability

Note: A GPU is built around an array of Streaming Multiprocessors (SMs) (see Hardware Implementation for more details). A multithreaded program is partitioned into blocks of threads that execute independently from each other, so that a GPU with more multiprocessors will automatically execute the program in less time than a GPU with fewer multiprocessors.

Kernels

CUDA c++ 通過允許程序員定義 c++ 函數(shù)( 稱為kernel )來擴(kuò)展 c++，當(dāng)調(diào)用這些函數(shù)時，由 N 個不同的 CUDA 線程并行執(zhí)行 N 次，而不是像常規(guī) c++ 函數(shù)那樣只執(zhí)行一次。

使用_ global _ 聲明說明符定義內(nèi)核，并使用新的<<<…>>>執(zhí)行配置語法(參見c++語言擴(kuò)展)。每個執(zhí)行內(nèi)核的線程都有一個惟一的線程ID，可以在內(nèi)核中通過內(nèi)置變量訪問該ID。

下面的示例代碼使用內(nèi)置變量 threadIdx ，將兩個大小為N的向量A和B相加，并將結(jié)果存儲到向量C中:

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

執(zhí)行VecAdd()的N個線程中的每一個都執(zhí)行一次成對的相加。

Thread Hierarchy

為了方便起見，threadIdx 是一個三分量的向量，因此可以使用一維、二維或三維線程索引來標(biāo)識線程，從而形成一維、二維或三維線程塊，稱為線程塊。這提供了一種很自然的方法來調(diào)用跨域元素(如向量、矩陣或體)的計(jì)算。

線程的索引和線程 ID 以一種直接的方式相互關(guān)聯(lián):

• 對于一維塊，它們是相同的
• 對于大小為(Dx, Dy)的二維塊，索引為(x, y)的線程ID為(x + y Dx);
• 對于大小為Dx, Dy, Dz的三維塊，索引為(x, y, z)的線程ID為(x + y Dx + z Dx Dy)。

例如，下面的代碼將兩個大小為NxN的矩陣A和B相加，并將結(jié)果存儲到矩陣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);
    ...
}

每個塊的線程數(shù)量是有限制的，因?yàn)橐粋€塊的所有線程都駐留在同一個流多處理器核心上，必須共享該核心的有限內(nèi)存資源。 在當(dāng)前的gpu上，一個線程塊可能包含多達(dá)1024個線程 。

但是，一個內(nèi)核可以由多個形狀相同的線程塊執(zhí)行，這樣 線程總數(shù)就等于每個塊的線程數(shù)乘以塊的數(shù)量 。

塊被組織成一維、二維或三維的線程塊網(wǎng)格，如圖所示。 網(wǎng)格中線程塊的數(shù)量通常由正在處理的數(shù)據(jù)的大小決定 ，數(shù)據(jù)的大小通常超過系統(tǒng)中處理器的數(shù)量。

圖2 Grid of Thread Blocks

每個塊的線程數(shù)和每個網(wǎng)格的塊數(shù)在<<<…>>>語法的類型可以是int或dim3。二維塊或網(wǎng)格可以像上面的例子中那樣指定。

網(wǎng)格中的每個塊都可以通過一個一維、二維或三維的惟一索引來標(biāo)識。該索引可以通過內(nèi)核中內(nèi)置的blockIdx變量訪問。 線程塊的維度可以在內(nèi)核中通過內(nèi)置的blockDim變量訪問 。

擴(kuò)展前面的MatAdd()示例以處理多個塊，代碼如下所示。

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

線程塊大小為16x16(256個線程)，雖然在本例中是任意的，但卻是常見的選擇 。網(wǎng)格是用足夠的塊創(chuàng)建的，每個矩陣元素都有一個線程。

為了簡單起見，本示例假設(shè)每個維度中每個網(wǎng)格的線程數(shù)可以被該維度中每個塊的線程數(shù)整除，盡管事實(shí)并非如此。

線程塊需要獨(dú)立執(zhí)行 :必須能夠以任何順序執(zhí)行它們，并行或串行。這種獨(dú)立性要求允許線程塊在任意數(shù)量的核上以任意順序調(diào)度，如圖1所示，這使程序員能夠編寫隨核數(shù)量擴(kuò)展的代碼。

塊中的線程可以通過共享內(nèi)存共享數(shù)據(jù)，并通過同步它們的執(zhí)行來協(xié)調(diào)內(nèi)存訪問，從而進(jìn)行協(xié)作 。更精確地說，可以通過調(diào)用__syncthreads()內(nèi)部函數(shù)來指定內(nèi)核中的同步點(diǎn);__syncthreads()充當(dāng)一個屏障，在允許任何線程繼續(xù)之前，塊中的所有線程都必須等待。除了__syncthreads()之外，Cooperative Groups API還提供了一組豐富的線程同步原語。

為了高效合作，共享內(nèi)存應(yīng)該是每個處理器核心附近的低延遲內(nèi)存(很像L1緩存)，并且__syncthreads()應(yīng)該是輕量級的。

Thread Block Clusters

隨著NVIDIA Compute Capability 9.0的引入 ，CUDA編程模型引入了一個可選的層次結(jié)構(gòu)級別，稱為 線程塊集群，它由線程塊組成 。 與線程塊中的線程被保證在流多處理器上同步調(diào)度類似，集群中的線程塊也被保證在GPU中的GPU處理集群(GPC)上同步調(diào)度 。

與線程塊類似，集群也被組織成一維、二維或三維，如圖3所示。一個集群中的線程塊數(shù)量可以由用戶定義， CUDA支持一個集群中最多8個線程塊作為可移植的集群大小 。線程塊集群大小是否超過8取決于體系結(jié)構(gòu)，可以使用cudaoccuancymaxpotentialclustersize API進(jìn)行查詢。

圖3 Grid of Thread Block Clusters

Note: In a kernel launched using cluster support, the gridDim variable still denotes the size in terms of number of thread blocks , for compatibility purposes. The rank of a block in a cluster can be found using the Cluster Group API.

線程塊集群可以在內(nèi)核中使用編譯器時間內(nèi)核屬性__cluster_dims__(X,Y,Z)或使用CUDA內(nèi)核啟動API cudaLaunchKernelEx來啟用。下面的示例展示了如何使用編譯器時間內(nèi)核屬性啟動集群。使用內(nèi)核屬性的集群大小在編譯時固定，然后可以使用經(jīng)典的<<<，>>>啟動內(nèi)核。如果內(nèi)核使用編譯時集群大小，則在啟動內(nèi)核時無法修改集群大小。

• 編譯期指定

// Kernel definition
// Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{

}

int main()
{
    float *input, *output;
    // Kernel invocation with compile time cluster size
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension is not affected by cluster launch, and is still enumerated
    // using number of blocks. 
    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<

• 運(yùn)行期指定

線程塊集群大小也可以在運(yùn)行時設(shè)置，并且可以使用CUDA內(nèi)核啟動API cudaLaunchKernelEx啟動內(nèi)核。下面的代碼示例展示了如何使用可擴(kuò)展API啟動集群內(nèi)核。

// Kernel definition
// No compile time attribute attached to the kernel
__global__ void cluster_kernel(float *input, float* output)
{
  
}

int main()
{
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    cluster_kernel<<

在具有9.0計(jì)算能力的GPU中，集群中的所有線程塊都被保證在單個GPU處理集群(GPC)上共同調(diào)度，并允許集群中的線程塊使用cluster Group API cluster.sync()執(zhí)行硬件支持的同步 。集群組還提供了成員函數(shù)，分別使用num_threads()和num_blocks() API根據(jù)線程數(shù)或塊數(shù)查詢集群組的大小?？梢苑謩e通過dim_threads()和dim_blocks() API查詢集群組中線程或塊的級別。

屬于一個集群的線程塊可以訪問 分布式共享內(nèi)存 。集群中的線程塊能夠?qū)Ψ植际焦蚕韮?nèi)存中的任何地址進(jìn)行讀、寫和執(zhí)行原子操作。分布式共享內(nèi)存給出了一個在分布式共享內(nèi)存中執(zhí)行直方圖的示例。