CUDA C++程序设计模型
CUDA C++程序設(shè)計模型
本章介紹了CUDA編程模型背后的主要概念，概述了它們在C++中的暴露方式。在編程接口中給出了CUDA C++的廣泛描述。
使用的矢量加法示例的完整代碼可以在矢量加法CUDA示例中找到。
一. 內(nèi)核
CUDA C++通過允許程序員定義C++函數(shù)，稱為內(nèi)核，擴展了C++，當(dāng)調(diào)用時，用n個不同的CUDA線程并行執(zhí)行n次，而不是像常規(guī)C++函數(shù)那樣只執(zhí)行一次。
一個內(nèi)核使用了__global__聲明說明符來定義，并且使用一個新的<<<…>執(zhí)行配置語法（參見C++語言擴展）來指定執(zhí)行給定內(nèi)核調(diào)用的內(nèi)核的CUDA線程的數(shù)量。執(zhí)行內(nèi)核的每個線程都有一個唯一的線程ID，可以通過內(nèi)置變量在內(nèi)核中訪問。
下面的示例代碼使用內(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);
…
}
二．線程層次結(jié)構(gòu)
為了方便起見，threadIdx是一個三分量向量，因此可以使用一維、二維或三維螺紋索引來識別螺紋，形成一維、二維或三維螺紋塊，稱為螺紋塊。這提供了一種跨域元素（如向量、矩陣或卷積）調(diào)用計算的自然方法。
一個線程的索引和它的線程ID以一種簡單的方式相互關(guān)聯(lián)：對于一維模塊，它們是相同的；對于二維模塊（Dx，Dy），索引（x，y）線程的線程ID是（x+ydx）；對于三維模塊（Dx，Dy，Dz），索引（x，y，z）線程的線程ID是（x+ydx+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īng)該位于同一個處理器核心上，并且必須共享該核心的有限內(nèi)存資源。在當(dāng)前GPU上，一個線程塊最多可以包含1024個線程。但是，內(nèi)核可以由多個形狀相同的線程塊執(zhí)行，因此線程總數(shù)等于每個塊的線程數(shù)乘以塊的數(shù)量。
塊被組織成一維、二維或三維的螺紋塊網(wǎng)格，如圖1所示。網(wǎng)格中線程塊的數(shù)量通常由正在處理的數(shù)據(jù)的大小決定，該大小通常超過系統(tǒng)中處理器的數(shù)量。
圖1. 螺紋塊網(wǎng)格
語法中指定的每個塊的線程數(shù)和每個網(wǎng)格的塊數(shù)可以是int或dim3類型。二維塊或網(wǎng)格可以在上面的示例中指定。
網(wǎng)格中的每個塊都可以由一個一維、二維或三維唯一索引標(biāo)識，該索引可通過內(nèi)置的blockIdx變量在內(nèi)核中訪問。線程塊的維度可以通過內(nèi)置的blockDim變量在內(nèi)核中訪問。擴展前面的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ù)均勻地整除，盡管事實并非如此。 線程塊需要獨立執(zhí)行：必須能夠以任何順序并行或串聯(lián)執(zhí)行它們。
這種獨立性要求允許線程塊在任意數(shù)量的內(nèi)核上按任意順序進行調(diào)度，使程序員能夠編寫隨內(nèi)核數(shù)量而擴展的代碼。
塊中的線程可以通過共享一些共享內(nèi)存來協(xié)作，并通過同步它們的執(zhí)行來協(xié)調(diào)內(nèi)存訪問。更準(zhǔn)確地說，可以通過調(diào)用syncthreads（）內(nèi)部函數(shù)來指定內(nèi)核中的同步點；syncthreads（）充當(dāng)一個屏障，塊中的所有線程都必須在該屏障上等待，然后才能允許任何線程繼續(xù)。共享內(nèi)存給出了一個使用共享內(nèi)存的示例。除了syncthreads（）之外，協(xié)作組API還提供了一組豐富的線程同步原語。
為了實現(xiàn)高效的協(xié)作，共享內(nèi)存應(yīng)該是靠近每個處理器核心的低延遲內(nèi)存（很像一級緩存），而syncthreads（）應(yīng)該是輕量級的。
三．內(nèi)存層次結(jié)構(gòu)
CUDA線程可以在執(zhí)行期間從多個內(nèi)存空間訪問數(shù)據(jù)，如圖2所示。每個線程都有專用的本地內(nèi)存。每個線程塊都有對該塊的所有線程可見的共享內(nèi)存，并且與該塊具有相同的生存期。所有線程都可以訪問相同的全局內(nèi)存。
所有線程還可以訪問另外兩個只讀內(nèi)存空間：常量和紋理內(nèi)存空間。全局、常量和紋理內(nèi)存空間針對不同的內(nèi)存使用進行了優(yōu)化（請參閱設(shè)備內(nèi)存訪問）。對于某些特定的數(shù)據(jù)格式（請參見紋理和曲面內(nèi)存），紋理內(nèi)存還提供不同的尋址模式以及數(shù)據(jù)過濾。
全局、常量和紋理內(nèi)存空間在同一個應(yīng)用程序啟動的內(nèi)核之間是持久的。
圖2. 內(nèi)存層次結(jié)構(gòu)
四．異構(gòu)程序設(shè)計
如圖3所示，CUDA編程模型假定CUDA線程在物理上獨立的設(shè)備上執(zhí)行，該設(shè)備作為運行C++程序的主機的協(xié)處理器。例如，當(dāng)內(nèi)核在GPU上執(zhí)行，而C++程序的其余部分在CPU上執(zhí)行時，情況就是這樣。
CUDA編程模型還假設(shè)主機和設(shè)備都在DRAM中保持各自獨立的內(nèi)存空間，分別稱為主機內(nèi)存和設(shè)備內(nèi)存。因此，程序通過調(diào)用CUDA運行時（在編程接口中描述）來管理內(nèi)核可見的全局、常量和紋理內(nèi)存空間。這包括設(shè)備內(nèi)存分配和釋放，以及主機和設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸。
統(tǒng)一內(nèi)存提供托管內(nèi)存以橋接主機和設(shè)備內(nèi)存空間。托管內(nèi)存可以從系統(tǒng)中的所有CPU和GPU訪問，作為具有公共地址空間的單個相干內(nèi)存映像。此功能支持設(shè)備內(nèi)存的超額訂閱，并且通過消除在主機和設(shè)備上顯式鏡像數(shù)據(jù)的需要，可以大大簡化移植應(yīng)用程序的任務(wù)。有關(guān)統(tǒng)一內(nèi)存的介紹，請參見統(tǒng)一內(nèi)存編程。
圖3. 異構(gòu)程序設(shè)計
注：串行代碼在主機上執(zhí)行，并行代碼在設(shè)備上執(zhí)行。
五．計算能力
設(shè)備的計算能力由版本號表示，有時也稱為“SM版本”。此版本號標(biāo)識GPU硬件支持的功能，并由應(yīng)用程序在運行時用于確定當(dāng)前GPU上可用的硬件功能和/或指令。計算能力包括主要修訂號X和次要修訂號Y，并用X.Y表示。
具有相同主要修訂號的設(shè)備具有相同的核心體系結(jié)構(gòu)。主要版本號為7（基于Volta架構(gòu)的設(shè)備）、6（基于Pascal架構(gòu)的設(shè)備）、5（基于Maxwell架構(gòu)的設(shè)備）、3（基于Kepler架構(gòu)的設(shè)備）、2（基于Fermi架構(gòu)的設(shè)備）和1（基于Tesla架構(gòu)的設(shè)備）。
次要修訂號對應(yīng)于核心架構(gòu)的增量改進，可能包括新功能。
Turing是計算能力為7.5的設(shè)備的架構(gòu)，是基于Volta架構(gòu)的增量更新。
啟用CUDA的GPU列出了所有啟用CUDA的設(shè)備及其計算能力。計算能力給出了每個計算能力的技術(shù)規(guī)范。
注：特定GPU的計算能力版本不應(yīng)與CUDA版本（如CUDA 7.5、CUDA 8、CUDA 9）混淆，后者是CUDA軟件平臺的版本。CUDA平臺被應(yīng)用程序開發(fā)人員用來創(chuàng)建運行在許多代GPU架構(gòu)上的應(yīng)用程序，包括尚未發(fā)明的未來GPU架構(gòu)。雖然CUDA平臺的新版本通常通過支持新GPU體系結(jié)構(gòu)的計算能力版本來增加對新GPU體系結(jié)構(gòu)的本機支持，但CUDA平臺的新版本通常還包括獨立于硬件生成的軟件功能。
從CUDA 7.0和CUDA 9.0開始，Tesla和Fermi架構(gòu)不再受支持。
總結(jié)
以上是生活随笔為你收集整理的CUDA C++程序设计模型的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： CUDA C++编程手册（总论）
- 下一篇： CUDA C++编程接口：编译