可伸縮的編程模型

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

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

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

圖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++ 函數( 稱為kernel )來擴展 c++，當調用這些函數時，由 N 個不同的 CUDA 線程并行執行 N 次，而不是像常規 c++ 函數那樣只執行一次。

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

下面的示例代碼使用內置變量 threadIdx ，將兩個大小為N的向量A和B相加，并將結果存儲到向量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);
    ...
}

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

Thread Hierarchy

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

線程的索引和線程 ID 以一種直接的方式相互關聯:

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

例如，下面的代碼將兩個大小為NxN的矩陣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);
    ...
}

每個塊的線程數量是有限制的，因為一個塊的所有線程都駐留在同一個流多處理器核心上，必須共享該核心的有限內存資源。 在當前的gpu上，一個線程塊可能包含多達1024個線程 。

但是，一個內核可以由多個形狀相同的線程塊執行，這樣 線程總數就等于每個塊的線程數乘以塊的數量 。

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

圖2 Grid of Thread Blocks

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

網格中的每個塊都可以通過一個一維、二維或三維的惟一索引來標識。該索引可以通過內核中內置的blockIdx變量訪問。 線程塊的維度可以在內核中通過內置的blockDim變量訪問 。

擴展前面的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個線程)，雖然在本例中是任意的，但卻是常見的選擇 。網格是用足夠的塊創建的，每個矩陣元素都有一個線程。

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

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

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

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

Thread Block Clusters

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

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

圖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.

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

• 編譯期指定

// 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<<

• 運行期指定

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

// 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計算能力的GPU中，集群中的所有線程塊都被保證在單個GPU處理集群(GPC)上共同調度，并允許集群中的線程塊使用cluster Group API cluster.sync()執行硬件支持的同步 。集群組還提供了成員函數，分別使用num_threads()和num_blocks() API根據線程數或塊數查詢集群組的大小。可以分別通過dim_threads()和dim_blocks() API查詢集群組中線程或塊的級別。

屬于一個集群的線程塊可以訪問 分布式共享內存 。集群中的線程塊能夠對分布式共享內存中的任何地址進行讀、寫和執行原子操作。分布式共享內存給出了一個在分布式共享內存中執行直方圖的示例。