本章和下一章中使用的向量加法示例的完整代碼可以在 vectorAddCUDA示例中找到。

2.1 內(nèi)核

CUDA C++ 通過(guò)允許程序員定義稱為kernel的 C++ 函數(shù)來(lái)擴(kuò)展 C++，當(dāng)調(diào)用內(nèi)核時(shí)，由 N 個(gè)不同的 CUDA 線程并行執(zhí)行 N 次，而不是像常規(guī) C++ 函數(shù)那樣只執(zhí)行一次。

使用__global__聲明說(shuō)明符定義內(nèi)核，并使用新的<<<...>>>執(zhí)行配置語(yǔ)法指定內(nèi)核調(diào)用的 CUDA 線程數(shù)（請(qǐng)參閱C++ 語(yǔ)言擴(kuò)展）。 每個(gè)執(zhí)行內(nèi)核的線程都有一個(gè)唯一的線程 ID，可以通過(guò)內(nèi)置變量在內(nèi)核中訪問(wèn)。

作為說(shuō)明，以下示例代碼使用內(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 分量向量，因此可以使用一維、二維或三維的線程索引來(lái)識(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ù)的大小決定，通常超過(guò)系統(tǒng)中的處理器數(shù)量。

<<<...>>>語(yǔ)法中指定的每個(gè)塊的線程數(shù)和每個(gè)網(wǎng)格的塊數(shù)可以是int或dim3類型。如上例所示，可以指定二維塊或網(wǎng)格。

網(wǎng)格中的每個(gè)塊都可以由一個(gè)一維、二維或三維的惟一索引標(biāo)識(shí)，該索引可以通過(guò)內(nèi)置的blockIdx變量在內(nèi)核中訪問(wèn)。線程塊的維度可以通過(guò)內(nèi)置的blockDim變量在內(nèi)核中訪問(wèn)。

擴(kuò)展前面的MatAdd()示例來(lái)處理多個(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è)線程來(lái)處理。為簡(jiǎn)單起見，本例假設(shè)每個(gè)維度中每個(gè)網(wǎng)格的線程數(shù)可以被該維度中每個(gè)塊的線程數(shù)整除，盡管事實(shí)并非如此。

程塊需要獨(dú)立執(zhí)行：必須可以以任何順序執(zhí)行它們，并行或串行。 這種獨(dú)立性要求允許跨任意數(shù)量的內(nèi)核以任意順序調(diào)度線程塊，如下圖所示，使程序員能夠編寫隨內(nèi)核數(shù)量擴(kuò)展的代碼。

塊內(nèi)的線程可以通過(guò)一些共享內(nèi)存共享數(shù)據(jù)并通過(guò)同步它們的執(zhí)行來(lái)協(xié)調(diào)內(nèi)存訪問(wèn)來(lái)進(jìn)行協(xié)作。 更準(zhǔn)確地說(shuō)，可以通過(guò)調(diào)用__syncthreads()內(nèi)部函數(shù)來(lái)指定內(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)存空間訪問(wèn)數(shù)據(jù)，如下圖所示。每個(gè)線程都有私有的本地內(nèi)存。 每個(gè)線程塊都具有對(duì)該塊的所有線程可見的共享內(nèi)存，并且具有與該塊相同的生命周期。 所有線程都可以訪問(wèn)相同的全局內(nèi)存。

還有兩個(gè)額外的只讀內(nèi)存空間可供所有線程訪問(wèn)：常量和紋理內(nèi)存空間。 全局、常量和紋理內(nèi)存空間針對(duì)不同的內(nèi)存使用進(jìn)行了優(yōu)化（請(qǐng)參閱設(shè)備內(nèi)存訪問(wèn)）。 紋理內(nèi)存還為某些特定數(shù)據(jù)格式提供不同的尋址模式以及數(shù)據(jù)過(guò)濾（請(qǐng)參閱紋理和表面內(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)存。因此，程序通過(guò)調(diào)用 CUDA 運(yùn)行時(shí)（在編程接口中描述）來(lái)管理內(nèi)核可見的全局、常量和紋理內(nèi)存空間。這包括設(shè)備內(nèi)存分配和釋放以及主機(jī)和設(shè)備內(nèi)存之間的數(shù)據(jù)傳輸。

統(tǒng)一內(nèi)存提供托管內(nèi)存來(lái)橋接主機(jī)和設(shè)備內(nèi)存空間。托管內(nèi)存可從系統(tǒng)中的所有 CPU 和 GPU 訪問(wèn)，作為具有公共地址空間的單個(gè)連貫內(nèi)存映像。此功能可實(shí)現(xiàn)設(shè)備內(nèi)存的超額訂閱，并且無(wú)需在主機(jī)和設(shè)備上顯式鏡像數(shù)據(jù)，從而大大簡(jiǎn)化了移植應(yīng)用程序的任務(wù)。有關(guān)統(tǒng)一內(nèi)存的介紹，請(qǐng)參閱統(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 編程模型通過(guò)異步編程模型為內(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ì)象來(lái)同步操作的完成。這樣的同步對(duì)象可以由用戶顯式管理（例如，cuda::memcpy_async）或在庫(kù)中隱式管理（例如，cooperative_groups::memcpy_async）。

同步對(duì)象可以是cuda::barrier或cuda::pipeline。這些對(duì)象在Asynchronous Barrier和Asynchronous Data Copies using cuda::pipeline.中進(jìn)行了詳細(xì)說(shuō)明。這些同步對(duì)象可以在不同的線程范圍內(nèi)使用。作用域定義了一組線程，這些線程可以使用同步對(duì)象與異步操作進(jìn)行同步。下表定義了CUDA c++中可用的線程作用域，以及可以與每個(gè)線程同步的線程。

這些線程作用域是在CUDA標(biāo)準(zhǔn)c++庫(kù)中作為標(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ū)以來(lái)，完成過(guò)上百場(chǎng)培訓(xùn)，幫助上萬(wàn)個(gè)開發(fā)者了解人工智能和 GPU 編程開發(fā)。在計(jì)算機(jī)視覺，高性能計(jì)算領(lǐng)域完成過(guò)多個(gè)獨(dú)立項(xiàng)目。并且，在機(jī)器人和無(wú)人機(jī)領(lǐng)域，有過(guò)豐富的研發(fā)經(jīng)驗(yàn)。對(duì)于圖像識(shí)別，目標(biāo)的檢測(cè)與跟蹤完成過(guò)多種解決方案。曾經(jīng)參與 GPU 版氣象模式GRAPES，是其主要研發(fā)者。

審核編輯：郭婷