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

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

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

圖3. 在GPU上執(zhí)行內(nèi)核。
CUDA為線程和塊定義了內(nèi)置的三維變量。線程使用內(nèi)置的三維變量threadIdx編制索引。三維索引提供了一種自然的方法來索引向量、矩陣和體積中的元素，并使CUDA編程更容易。類似地，塊也使用名為blockIdx的內(nèi)置三維變量編制索引。
以下是幾個值得注意的要點：
CUDA架構(gòu)限制每個塊的線程數(shù)（每個塊限制1024個線程）。
線程塊的維度可以通過內(nèi)置的blockDim變量在內(nèi)核中訪問。
syncu中的線程可以使用syncu函數(shù)同步。使用同步線程時，塊中的所有線程都必須等待，然后才能繼續(xù)。
在<<…>>>語法中指定的每個塊的線程數(shù)和每個網(wǎng)格的塊數(shù)可以是int或dim3類型。這些三角括號標(biāo)記從主機代碼到設(shè)備代碼的調(diào)用。它也被稱為內(nèi)核啟動。
下面用于添加兩個矩陣的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有一個內(nèi)存層次結(jié)構(gòu)，如圖4所示。

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

總結(jié)

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

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