CUDA 技術手冊翻譯 Day 2

「CUDA 技術手冊翻譯」這一系列文章是個人的嘗試翻譯，原文是 CUDA C Programming Giude 文檔 ，而標題的 Day N 分篇則源自於文章底部的原文網址。
若有任何錯誤還煩請各位指出，希望能在翻譯與閱讀的同時增進對 Nvidia CUDA 的認識。個人也會在進一步的學習後，於文章末分享自己的心得。

2. 編程模型

• 此章節介紹CUDA編程模型背後的主要概念，概述如何使用C表達他們
• 所有在此章節與次個章節使用的向量加法範例，皆可以在 CUDA 範例中 vectorAdd 裡找到。

2.1 內核

CUDA C 允許程序員藉由定義 C 語言的方法與呼叫核心，延伸 C 語言，當被呼叫時，被 N 個不同的 CUDA 線程並行執行 N 次，而不是像普通的 C 函數一樣只有一次。

一個核心使用 __global__ 聲明說明符定義，而對於特定的內核調用，執行該內核的CUDA線程的數量是使用新的 <<< ... >>> 執行配置語法。每一個執行內核的線程被賦予一個獨一無二的線程ID，以此通過內置的 threadIdx 變量訪問內核。

如下插圖，下面的示例代碼添加兩個大小為 N 的向量 A 和 B ，並將結果存儲到向量 C 中：

在這裡，執行 VecAdd（） 的 N 個線程都執行一次一對一的加法(兩兩相加)

2.2 線程層次結構

為了方便， threadIdx 是一個三維向量，所以可以用一維、二維、三維下標來指定唯一線程。形成一維、二維或三維的線程區塊，稱為線程塊。這提供了一種自然的方式來調用跨越元素的計算，例如向量，矩陣或體積。

一個線程的索引和它的線程ID以直接的方式相互關聯，對於一維塊，它們是相同的﹔對於尺寸為（Dx，Dy）的二維塊，索引（x，y）的線程的線程ID為（x + y Dx）﹔對於尺寸為（Dx，Dy，Dz）的三維塊，索引（x，y，z）的線程的線程ID為（x + y Dx + z Dx Dy）。

如下範例，下面的代碼相加了兩個大小為 NxN 的矩陣 A 和 B ，並將結果存儲到矩陣 C 中:

每塊的線程數量是有限制的，因為一個塊的所有線程都應該駐留在同一個處理器內核上，並且必須共享該內核的有限內存資源。在當前的GPU上，一個線程塊最多可以包含1024個線程。

但是，一個內核可以由多個相同形狀的線程塊執行，因此線程的總數等於每塊的線程數乘以塊的數量。

如圖6所示，區塊被一維，二維或三維的線程塊組織成網格。網格中線程塊的數量，通常會由數據程序的大小或系統的程序數量來決定，因此它可以大大超過限制。

Figure 6. Grid of Thread Blocks

每個塊的線程數和<<< … >>>語法中指定的每個網格的塊數可以是int類型或dim3類型。可以像上面的例子那樣指定二維塊或網格。

網格中的每個塊都可以通過內核中通過內置blockIdx變量在內核中訪問的一維，二維或三維索引來標識。線程塊的維度可以通過內置的blockDim變量在內核中訪問。

擴展前面的MatAdd（）示例以處理多個塊，代碼變為如下所示。

儘管在這種情況下，我們可以任意的選擇線程塊大小，但16x16（256線程）是常見的選擇。如同以前每個矩陣元素都有一個線程，每個網格需要足夠的區塊來創造。為了簡單起見，本示例假定每個維度中每個網格的線程數可以被該維度中每個塊的線程數整除，但現實中並不一定如此。

線程塊需要獨立執行：它必須能夠以任意順序，並行或串行執行。這種獨立性要求允許按照圖5所示的任意數量的內核以任意順序調度線程塊，從而使編程人員能夠編寫與內核數量一致的代碼。

塊內的線程可以通過共享內存共享數據並通過同步執行來協調內存訪問來進行協作。更準確地說，可以通過調用 __syncthreads（） 內部函數來指定內核中的同步點;__syncthreads（）充當一個分界點，在允許繼續進行之前，塊中的所有線程必須等待。共享內存提供了一個使用共享內存的例子。除 __syncthreads（） 外，Cooperative Groups API還提供了豐富的線程同步原語。

為了有效協作，共享內存預計將是每個處理器內核附近的低延遲內存（類似於L1緩存），並且__syncthreads（）預計會很輕量。

2.3 內存層次結構

CUDA線程可能會允許在執行期間從多個內存空間訪問數據，如圖7所示。每個線程都有私有本地內存。每個線程塊都具有對該塊的所有線程都可見的共享存儲器，並且具有與該塊相同的生命週期。所有線程都可以訪問相同的全局內存。

所有線程還可以訪問另外兩個只讀內存空間：常量和紋理內存空間。全局，常量和紋理內存空間針對不同的內存使用進行了優化。對於某些特定的數據格式，紋理內存還提供了不同的尋址模式以及數據過濾。

全局，常量和紋理內存空間在同一應用程序的內核啟動之間是連續的。
圖7.內存層次結構

2.4 異構编程

如圖8所示，CUDA編程模型假定CUDA線程在物理上分離的設備上執行，該設備作為運行C程序的主機的協處理器運行。例如，當內核在GPU上執行時，C程序的其餘部分在CPU上執行。

CUDA編程模型還假定主機和設備都在DRAM中分別維護其自己的獨立存儲空間，分別稱為主機存儲器和設備存儲器。因此，程序通過調用CUDA運行庫來管理內核可見的全局，常量和紋理內存空間。這包括設備內存分配和重新分配以及主機和設備內存之間的數據傳輸。

統一內存提供託管內存以橋接主機和設備內存空間。託管內存可以從系統中的所有CPU和GPU訪問，作為具有公共地址空間的單連貫內存映像。這種功能可以實現對設備內存的超額訂購，並且可以通過消除在主機和設備上顯式鏡像數據的需求，大大簡化移植應用程序的任務。

圖8.異構編程

附註：串行代碼在主機上執行，而並行代碼在設備上執行。

2.5 計算能力

設備的計算能力由版本號表示，有時也稱為“SM版本”。此版本號標識GPU硬件支持的功能，並由運行時的應用程序用於確定當前GPU上可用的硬件功能和/或指令。

計算能力包括主版本號X和次版本號Y，用X.Y表示。

具有相同主要版本號的設備具有相同的核心架構。基於Volta架構的設備的主要版本號為7，基於Pascal架構的設備為6，基於Maxwell架構的設備為5，基於Kepler架構的設備為3，基於Fermi架構的設備為2，和基於Teslaarchitecture的設備為1。

次要修訂號對應於核心架構的增量改進，可能包括新功能。計算能力給出了每種計算能力的技術規格。

注意：不應該將特定GPU的計算能力版本與CUDA軟件平台版本的CUDA版本（例如，CUDA 7.5，CUDA 8，CUDA 9）混淆。應用程序開發人員使用CUDA平台創建在多代GPU架構上運行的應用程序，包括尚未發明的未來GPU架構。雖然CUDA平台的新版本通常通過支持該體系結構的計算功能版本為新的GPU架構添加本地支持，但新版本的CUDA平台通常還包含獨立於硬件生成的軟件功能。

附註：從CUDA 7.0和CUDA 9.0開始，不再支持Tesla和Fermi架構。

本文备注/经验分享：
each of the N threads that execute VecAdd() performs one pair-wise addition
整体翻译的话，可以翻译为“每个线程进行一对数值的加法”，请注意pair-wise addition还有另外一个意思是log2(N)方式的相加。如果你有16个浮点数，一种并行化的方式是：分成前8个，和后8个。前8个里面分成4个+4个，4个+4个分成2+2+2+2…这种pair-wise的累加是为了保持精度。 我们常见的，常说的shared memory上的规约，实际上就是这种累加。 所以也叫log2规约加法。 这种累加能增加精度，减少误差。 回到VectorAdd这个例子，这里面就是普通的两个数相加的意思。

a kernel can be executed by multiple equally-shaped thread blocks, 这里equally-shape是相同形状，因为我们启动kernel的时候，标准的runtime api语法是：<<<A,B>>>，这代表启动A个blocks，每个blocks都是B个线程的形状。

所以这里提到了：equally shaped。 CUDA也不支持一次启动中有多种不同形状的block，如果需要有多种不同形状的blocks，可以多次启动，或者自己用代码变形。

A thread block size of 16x16 (256 threads), although arbitrary in this case。这里arbitrary是任意的意思。我们启动kernel的时候，可以使用任意形状。但128， 256（本例), 512这些，是常见的形状选择。（16 * 16 = 256)。需要说明的是，某些kernel往往有个最佳的block形状，此形状下启动性能最好。但不能提前知道是什么样子，得反复试验，无直接的公式能告诉大家什么形状是最好的。注意很多时候我们选择形状的时候，需要加上if限制，if或者while (…)。 这是很多情况下问题规模并不能直接被你刚才选择的“同样形状/大小”的block给整除。此时往往需要过多启动blocks，并同时用if限制掉越界的线程。 这同时也是因为我们刚才说的，“同样形状”导致的----边界上的blocks并不能选择一些较小的，不同形状的。 注意OpenCL有不同的选择，它允许边界处的groups（等于CUDA的blocks）具有不同的形状，但这额外的增加了kernel书写者面对的复杂性。 CUDA比较易用，直接不让你考虑这样的。所以我们需要if或者while或者for之类的设定条件，处理好边界。