本章和下一章中使用的向量加法示例的完整代碼可以在 vectorAddCUDA示例中找到。
2.1 內(nèi)核
CUDA C++ 通過允許程序員定義稱為kernel的 C++ 函數(shù)來擴(kuò)展 C++,當(dāng)調(diào)用內(nèi)核時(shí),由 N 個(gè)不同的 CUDA 線程并行執(zhí)行 N 次,而不是像常規(guī) C++ 函數(shù)那樣只執(zhí)行一次。
使用__global__聲明說明符定義內(nèi)核,并使用新的<<<...>>>執(zhí)行配置語法指定內(nèi)核調(diào)用的 CUDA 線程數(shù)(請參閱C++ 語言擴(kuò)展)。 每個(gè)執(zhí)行內(nèi)核的線程都有一個(gè)唯一的線程 ID,可以通過內(nèi)置變量在內(nèi)核中訪問。
作為說明,以下示例代碼使用內(nèi)置變量threadIdx將兩個(gè)大小為 N 的向量 A 和 B 相加,并將結(jié)果存儲(chǔ)到向量 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 個(gè)線程中的每一個(gè)線程都會(huì)執(zhí)行一個(gè)加法。
2.2 線程層次
為方便起見,threadIdx 是一個(gè) 3 分量向量,因此可以使用一維、二維或三維的線程索引來識(shí)別線程,形成一個(gè)一維、二維或三維的線程塊,稱為block。 這提供了一種跨域的元素(例如向量、矩陣或體積)調(diào)用計(jì)算的方法。
線程的索引和它的線程 ID 以一種直接的方式相互關(guān)聯(lián):對(duì)于一維塊,它們是相同的; 對(duì)于大小為(Dx, Dy)的二維塊,索引為(x, y)的線程的線程ID為(x + y*Dx); 對(duì)于大小為 (Dx, Dy, Dz) 的三維塊,索引為 (x, y, z) 的線程的線程 ID 為 (x + y*Dx + z*Dx*Dy)。
例如,下面的代碼將兩個(gè)大小為NxN的矩陣A和B相加,并將結(jié)果存儲(chǔ)到矩陣C中:
// 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<<>>(A, B, C);
...
}
每個(gè)塊的線程數(shù)量是有限制的,因?yàn)橐粋€(gè)塊的所有線程都應(yīng)該駐留在同一個(gè)處理器核心上,并且必須共享該核心有限的內(nèi)存資源。在當(dāng)前的gpu上,一個(gè)線程塊可能包含多達(dá)1024個(gè)線程。
但是,一個(gè)內(nèi)核可以由多個(gè)形狀相同的線程塊執(zhí)行,因此線程總數(shù)等于每個(gè)塊的線程數(shù)乘以塊數(shù)。
塊被組織成一維、二維或三維的線程塊網(wǎng)格(grid),如下圖所示。網(wǎng)格中的線程塊數(shù)量通常由正在處理的數(shù)據(jù)的大小決定,通常超過系統(tǒng)中的處理器數(shù)量。
<<<...>>>語法中指定的每個(gè)塊的線程數(shù)和每個(gè)網(wǎng)格的塊數(shù)可以是int或dim3類型。如上例所示,可以指定二維塊或網(wǎng)格。
網(wǎng)格中的每個(gè)塊都可以由一個(gè)一維、二維或三維的惟一索引標(biāo)識(shí),該索引可以通過內(nèi)置的blockIdx變量在內(nèi)核中訪問。線程塊的維度可以通過內(nèi)置的blockDim變量在內(nèi)核中訪問。
擴(kuò)展前面的MatAdd()示例來處理多個(gè)塊,代碼如下所示。
// 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<<>>(A, B, C);
...
}
線程塊大小為16×16(256個(gè)線程),盡管在本例中是任意更改的,但這是一種常見的選擇。網(wǎng)格是用足夠的塊創(chuàng)建的,這樣每個(gè)矩陣元素就有一個(gè)線程來處理。為簡單起見,本例假設(shè)每個(gè)維度中每個(gè)網(wǎng)格的線程數(shù)可以被該維度中每個(gè)塊的線程數(shù)整除,盡管事實(shí)并非如此。
程塊需要獨(dú)立執(zhí)行:必須可以以任何順序執(zhí)行它們,并行或串行。 這種獨(dú)立性要求允許跨任意數(shù)量的內(nèi)核以任意順序調(diào)度線程塊,如下圖所示,使程序員能夠編寫隨內(nèi)核數(shù)量擴(kuò)展的代碼。
塊內(nèi)的線程可以通過一些共享內(nèi)存共享數(shù)據(jù)并通過同步它們的執(zhí)行來協(xié)調(diào)內(nèi)存訪問來進(jìn)行協(xié)作。 更準(zhǔn)確地說,可以通過調(diào)用__syncthreads()內(nèi)部函數(shù)來指定內(nèi)核中的同步點(diǎn);__syncthreads()充當(dāng)屏障,塊中的所有線程必須等待,然后才能繼續(xù)。Shared Memory給出了一個(gè)使用共享內(nèi)存的例子。 除了__syncthreads()之外,Cooperative Groups API還提供了一組豐富的線程同步示例。
為了高效協(xié)作,共享內(nèi)存是每個(gè)處理器內(nèi)核附近的低延遲內(nèi)存(很像 L1 緩存),并且__syncthreads()是輕量級(jí)的。
2.3 存儲(chǔ)單元層次
CUDA 線程可以在執(zhí)行期間從多個(gè)內(nèi)存空間訪問數(shù)據(jù),如下圖所示。每個(gè)線程都有私有的本地內(nèi)存。 每個(gè)線程塊都具有對(duì)該塊的所有線程可見的共享內(nèi)存,并且具有與該塊相同的生命周期。 所有線程都可以訪問相同的全局內(nèi)存。
還有兩個(gè)額外的只讀內(nèi)存空間可供所有線程訪問:常量和紋理內(nèi)存空間。 全局、常量和紋理內(nèi)存空間針對(duì)不同的內(nèi)存使用進(jìn)行了優(yōu)化(請參閱設(shè)備內(nèi)存訪問)。 紋理內(nèi)存還為某些特定數(shù)據(jù)格式提供不同的尋址模式以及數(shù)據(jù)過濾(請參閱紋理和表面內(nèi)存)。
全局、常量和紋理內(nèi)存空間在同一應(yīng)用程序的內(nèi)核啟動(dòng)中是持久的。
2.4 異構(gòu)編程
如下圖所示,CUDA 編程模型假定 CUDA 線程在物理獨(dú)立的設(shè)備上執(zhí)行,該設(shè)備作為運(yùn)行 C++ 程序的主機(jī)的協(xié)處理器運(yùn)行。例如,當(dāng)內(nèi)核在 GPU 上執(zhí)行而 C++ 程序的其余部分在 CPU 上執(zhí)行時(shí),就是這種情況。
CUDA 編程模型還假設(shè)主機(jī)(host)和設(shè)備(device)都在 DRAM 中維護(hù)自己獨(dú)立的內(nèi)存空間,分別稱為主機(jī)內(nèi)存和設(shè)備內(nèi)存。因此,程序通過調(diào)用 CUDA 運(yùn)行時(shí)(在編程接口中描述)來管理內(nèi)核可見的全局、常量和紋理內(nèi)存空間。這包括設(shè)備內(nèi)存分配和釋放以及主機(jī)和設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸。
統(tǒng)一內(nèi)存提供托管內(nèi)存來橋接主機(jī)和設(shè)備內(nèi)存空間。托管內(nèi)存可從系統(tǒng)中的所有 CPU 和 GPU 訪問,作為具有公共地址空間的單個(gè)連貫內(nèi)存映像。此功能可實(shí)現(xiàn)設(shè)備內(nèi)存的超額訂閱,并且無需在主機(jī)和設(shè)備上顯式鏡像數(shù)據(jù),從而大大簡化了移植應(yīng)用程序的任務(wù)。有關(guān)統(tǒng)一內(nèi)存的介紹,請參閱統(tǒng)一內(nèi)存編程。
注:串行代碼在主機(jī)(host)上執(zhí)行,并行代碼在設(shè)備(device)上執(zhí)行。
2.5 異步SIMT編程模型
在 CUDA 編程模型中,線程是進(jìn)行計(jì)算或內(nèi)存操作的最低抽象級(jí)別。 從基于 NVIDIA Ampere GPU 架構(gòu)的設(shè)備開始,CUDA 編程模型通過異步編程模型為內(nèi)存操作提供加速。 異步編程模型定義了與 CUDA 線程相關(guān)的異步操作的行為。
異步編程模型為 CUDA 線程之間的同步定義了異步屏障的行為。 該模型還解釋并定義了如何使用 cuda::memcpy_async 在 GPU計(jì)算時(shí)從全局內(nèi)存中異步移動(dòng)數(shù)據(jù)。
2.5.1 異步操作
異步操作定義為由CUDA線程發(fā)起的操作,并且與其他線程一樣異步執(zhí)行。在結(jié)構(gòu)良好的程序中,一個(gè)或多個(gè)CUDA線程與異步操作同步。發(fā)起異步操作的CUDA線程不需要在同步線程中.
這樣的異步線程(as-if 線程)總是與發(fā)起異步操作的 CUDA 線程相關(guān)聯(lián)。異步操作使用同步對(duì)象來同步操作的完成。這樣的同步對(duì)象可以由用戶顯式管理(例如,cuda::memcpy_async)或在庫中隱式管理(例如,cooperative_groups::memcpy_async)。
同步對(duì)象可以是cuda::barrier或cuda::pipeline。這些對(duì)象在Asynchronous Barrier和Asynchronous Data Copies using cuda::pipeline.中進(jìn)行了詳細(xì)說明。這些同步對(duì)象可以在不同的線程范圍內(nèi)使用。作用域定義了一組線程,這些線程可以使用同步對(duì)象與異步操作進(jìn)行同步。下表定義了CUDA c++中可用的線程作用域,以及可以與每個(gè)線程同步的線程。
這些線程作用域是在CUDA標(biāo)準(zhǔn)c++庫中作為標(biāo)準(zhǔn)c++的擴(kuò)展實(shí)現(xiàn)的。
2.6 Compute Capability
設(shè)備的Compute Capability由版本號(hào)表示,有時(shí)也稱其“SM版本”。該版本號(hào)標(biāo)識(shí)GPU硬件支持的特性,并由應(yīng)用程序在運(yùn)行時(shí)使用,以確定當(dāng)前GPU上可用的硬件特性和指令。
Compute Capability包括一個(gè)主要版本號(hào)X和一個(gè)次要版本號(hào)Y,用X.Y表示
主版本號(hào)相同的設(shè)備具有相同的核心架構(gòu)。設(shè)備的主要修訂號(hào)是8,為NVIDIA Ampere GPU的體系結(jié)構(gòu)的基礎(chǔ)上,7基于Volta設(shè)備架構(gòu),6設(shè)備基于Pascal架構(gòu),5設(shè)備基于Maxwell架構(gòu),3基于Kepler架構(gòu)的設(shè)備,2設(shè)備基于Fermi架構(gòu),1是基于Tesla架構(gòu)的設(shè)備。
次要修訂號(hào)對(duì)應(yīng)于對(duì)核心架構(gòu)的增量改進(jìn),可能包括新特性。
Turing是計(jì)算能力7.5的設(shè)備架構(gòu),是基于Volta架構(gòu)的增量更新。
CUDA-Enabled GPUs列出了所有支持 CUDA 的設(shè)備及其計(jì)算能力。Compute Capabilities給出了每個(gè)計(jì)算能力的技術(shù)規(guī)格。
關(guān)于作者
Ken He 是 NVIDIA 企業(yè)級(jí)開發(fā)者社區(qū)經(jīng)理 & 高級(jí)講師,擁有多年的 GPU 和人工智能開發(fā)經(jīng)驗(yàn)。自 2017 年加入 NVIDIA 開發(fā)者社區(qū)以來,完成過上百場培訓(xùn),幫助上萬個(gè)開發(fā)者了解人工智能和 GPU 編程開發(fā)。在計(jì)算機(jī)視覺,高性能計(jì)算領(lǐng)域完成過多個(gè)獨(dú)立項(xiàng)目。并且,在機(jī)器人和無人機(jī)領(lǐng)域,有過豐富的研發(fā)經(jīng)驗(yàn)。對(duì)于圖像識(shí)別,目標(biāo)的檢測與跟蹤完成過多種解決方案。曾經(jīng)參與 GPU 版氣象模式GRAPES,是其主要研發(fā)者。
審核編輯:郭婷
-
NVIDIA
+關(guān)注
關(guān)注
14文章
4814瀏覽量
102630 -
gpu
+關(guān)注
關(guān)注
27文章
4631瀏覽量
128440 -
人工智能
+關(guān)注
關(guān)注
1789文章
46316瀏覽量
236469
發(fā)布評(píng)論請先 登錄
相關(guān)推薦
評(píng)論