CUDA刷新器：CUDA編程模型
CUDA Refresher: The CUDA Programming Model
CUDA，CUDA刷新器，并行編程
這是CUDA更新系列的第四篇文章，它的目標是刷新CUDA中的關鍵概念、工具和初級或中級開發(fā)人員的優(yōu)化。
CUDA編程模型提供了GPU體系結構的抽象，它充當了應用程序與其在GPU硬件上的可能實現(xiàn)之間的橋梁。這篇文章概述了CUDA編程模型的主要概念，概述了它如何在通用編程語言如C/C++中暴露出來。
介紹一下CUDA編程模型中常用的兩個關鍵詞：主機和設備。
主機是系統(tǒng)中可用的CPU。與CPU相關聯(lián)的系統(tǒng)內存稱為主機內存。GPU被稱為設備，GPU內存也被稱為設備內存。
要執(zhí)行任何CUDA程序，有三個主要步驟：
將輸入數(shù)據(jù)從主機內存復制到設備內存，也稱為主機到設備傳輸。
加載GPU程序并執(zhí)行，在片上緩存數(shù)據(jù)以提高性能。
將結果從設備內存復制到主機內存，也稱為設備到主機傳輸。
CUDA內核和線程層次結構
圖1顯示了CUDA內核是一個在GPU上執(zhí)行的函數(shù)。應用程序的并行部分由K個不同的CUDA線程并行執(zhí)行k次，而不是像常規(guī)C/C++函數(shù)那樣只進行一次。

Figure 1. The kernel is a function executed on the GPU.
每一個CUDA內核都以一個__global__聲明說明符開頭。程序員通過使用內置變量為每個線程提供唯一的全局ID。

圖2. CUDA內核被細分為塊。
一組線程稱為CUDA塊。CUDA塊被分組到一個網(wǎng)格中。內核作為線程塊的網(wǎng)格來執(zhí)行（圖2）。
每個CUDA塊由一個流式多處理器（SM）執(zhí)行，不能遷移到GPU中的其他SMs（搶占、調試或CUDA動態(tài)并行期間除外）。一個SM可以根據(jù)CUDA塊所需的資源運行多個并發(fā)CUDA塊。每個內核在一個設備上執(zhí)行，CUDA支持一次在一個設備上運行多個內核。圖3顯示了GPU中可用硬件資源的內核執(zhí)行和映射。

圖3. 在GPU上執(zhí)行內核。
CUDA為線程和塊定義了內置的三維變量。線程使用內置的三維變量threadIdx編制索引。三維索引提供了一種自然的方法來索引向量、矩陣和體積中的元素，并使CUDA編程更容易。類似地，塊也使用名為blockIdx的內置三維變量編制索引。
以下是幾個值得注意的要點：
CUDA架構限制每個塊的線程數(shù)（每個塊限制1024個線程）。
線程塊的維度可以通過內置的blockDim變量在內核中訪問。
syncu中的線程可以使用syncu函數(shù)同步。使用同步線程時，塊中的所有線程都必須等待，然后才能繼續(xù)。
在<<…>>>語法中指定的每個塊的線程數(shù)和每個網(wǎng)格的塊數(shù)可以是int或dim3類型。這些三角括號標記從主機代碼到設備代碼的調用。它也被稱為內核啟動。
下面用于添加兩個矩陣的CUDA程序顯示多維blockIdx和threadIdx以及blockDim等其他變量。在下面的例子中，為了便于索引，選擇了一個2D塊，每個塊有256個線程，x和y方向各有16個線程。使用數(shù)據(jù)大小除以每個塊的大小來計算塊的總數(shù)。
// Kernel - Adding two matrices MatA and MatB
global void MatAdd(float MatA[N][N], float MatB[N][N],
float MatC[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
MatC[i][j] = MatA[i][j] + MatB[i][j];
}

int main()
{
…
// Matrix addition kernel launch from host code
dim3 threadsPerBlock(16, 16);
dim3 numBlocks((N + threadsPerBlock.x -1) / threadsPerBlock.x, (N+threadsPerBlock.y -1) / threadsPerBlock.y);
MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC);
…
}
Memory hierarchy
支持CUDA的GPU有一個內存層次結構，如圖4所示。

圖4. gpu中的內存層次結構。
以下內存由GPU架構公開：
這些寄存器對每個線程都是私有的，這意味著分配給線程的寄存器對其他線程不可見。編譯器決定寄存器的利用率。
一級/共享內存（SMEM）-每個SM都有一個快速的片上草稿行內存，可用作一級緩存和共享內存。CUDA塊中的所有線程都可以共享共享內存，在給定SM上運行的所有CUDA塊都可以共享SM提供的物理內存資源。。
只讀內存每個SM都有一個指令緩存、常量內存、紋理內存和對內核代碼只讀的RO緩存。
二級緩存二級緩存在所有SMs中共享，因此每個CUDA塊中的每個線程都可以訪問該內存。nvidiaa100 GPU已經(jīng)將二級緩存大小增加到40mb，而v100gpu中只有6mb。
全局內存這是位于GPU中的GPU和DRAM的幀緩沖區(qū)大小。
NVIDIA CUDA編譯器在優(yōu)化內存資源方面做得很好，但專家CUDA開發(fā)人員可以選擇有效地使用這種內存層次結構來優(yōu)化CUDA程序。
計算能力
GPU的計算能力決定了GPU硬件支持的通用規(guī)范和可用特性。此版本號可由應用程序在運行時使用，以確定當前GPU上可用的硬件功能或指令。
每個GPU都有一個版本號，表示為X.Y，其中X包括主要修訂號，Y包含次要修訂號。小版本號對應于架構的增量改進，可能包括新特性。
有關任何支持CUDA的設備的計算能力的更多信息，請參閱CUDA示例代碼設備查詢。此示例枚舉系統(tǒng)中存在的CUDA設備的屬性
摘要
CUDA編程模型提供了一種異構環(huán)境，其中主機代碼在CPU上運行C/C++程序，內核在物理上分離的GPU設備上運行。CUDA編程模型還假設主機和設備都保持各自獨立的內存空間，分別稱為主機內存和設備內存。CUDA代碼還通過PCIe總線提供主機和設備內存之間的數(shù)據(jù)傳輸。
CUDA還公開了許多內置變量，并提供了多維索引的靈活性，以簡化編程。CUDA還管理不同的內存，包括寄存器、共享內存和一級緩存、二級緩存和全局內存。高級開發(fā)人員可以有效地使用這些內存來優(yōu)化CUDA程序。

總結

以上是生活随笔為你收集整理的CUDA刷新器：CUDA编程模型的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內容還不錯，歡迎將生活随笔推薦給好友。