CUDA C程式設計手冊: 程式設計模型

• 核函數
• 線程體系
• 記憶體體系
• 異構程式設計

核函數

CUDA C 是對C的一中擴充， 允許程式員定義自己的C函數， 并且将之稱為 核。與傳統C函數不同的是， 這樣的核被調用的時候， 會被CUDA 線程并行地執行N次。

核函數定義的時候使用 _global_聲明辨別符。同時， 在調用這樣的核函數的時候， 會使用一個特定的文法

<<< … >>>來指定執行的配置。每一個線程執行這個核函數的時候， 都會被配置設定一個獨一無二的線程ID， 這個線程ID可以通過内置的變量 threadIdx來獲得。

下面是一個示例代碼，示範兩個大小為N的向量A和B相加，結果存放在C中。 其中每個線程都會執行VecAdd()核函數， 使得每一對元素相加。

// 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);
	...
}
           

線程體系

為了簡單起見， threadIdx是一個三元素的向量， 是以可以通過一個一維/二維/三維的線程索引來組成一個一維/二維/三維的線程塊(thread block)。對于向量、矩陣和體資料的計算來說， 這是一種非常自然且合适的方法。

線程的索引和其ID之間的關系是非常直覺的：對于一維的線程塊， 他們是一樣的， 即索引就是線程的ID； 對于大小為( D x D_x Dx​, D y ） D_y） Dy​）的線程塊來說， 線程索引為 ( x , y ) (x, y) (x,y)的ID是 ( x + y ⋅ D x ) (x+y \cdot D_x) (x+y⋅Dx​); 對于一個三維 ( D x , D y , D z ) (D_x, D_y, D_z) (Dx​,Dy​,Dz​)的線程塊來說， 線程的索引為 ( x , y , z ) (x, y, z) (x,y,z)的ID是 ( x + y ⋅ D x + z ⋅ D x ⋅ D z ) (x+y\cdot D_x+ z\cdot D_x \cdot D_z) (x+y⋅Dx​+z⋅Dx​⋅Dz​)。

下面簡單的示例代碼，表示的是兩個 N × N N \times N N×N的矩陣A和B相加，結果存儲在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);
	...
}
           

每個線程塊中的線程數量都有上限，因為每個線程塊中的線程都會駐留在同一個處理核心中且同時共享有限的處理核心上的記憶體資源。對于現在發行的GPU， 一個線程塊中線程的上限是1024。然而， 一個核函數， 一個核函數可以被多個線程塊執行， 是以線程的總數量為每個線程塊中線程數量與線程塊個數的乘積。

線程塊可以被組成一維的， 二維的或者三維的線程塊網格grid。線程塊網格中線程塊的數量通常根據需要處理資料大小而得來， 或者是系統的處理器的數量， 這樣能獲得較大的加速效果。

線程塊的執行是無序的， 即他們的執行時互相獨立的， 沒有依賴關系。同一個線程塊内的線程可以通過共享記憶體shared memory來進行合作， 或者通過線程同步來“有序”的通路記憶體。更确切地說， 可以通過在核函數中調用*__syncthreads()*來指明的同步點， 這個函數類似于一個栅欄， 隻有當塊中的所有線程都到達了該點才可以繼續往下執行，否則先到達該點的線程都需要停下來等待其他線程。

記憶體體系

CUDA中的線程在執行的過程中， 可以通路多種記憶體空間中的資料。每個線程都有自己的私有的本地記憶體。 每個線程塊都有對所有線程塊可見的共享記憶體且該記憶體與線程塊具有同樣的生命周期；所有的線程都可以通路同樣的全局記憶體空間。

存在兩種可以被所有線程通路的額外的隻讀記憶體空間：常量記憶體空間和紋理記憶體空間。全局、常量和紋理記憶體空間都存在一些優化的使用方法。紋理記憶體也提供了一些特殊的尋址模式和對一些特殊資料格式的資料濾波。全局、常量和紋理記憶體空間貫穿于整個應用。

異構程式設計

CUDA程式設計模型假定在實體分離的裝置device上執行的CUDA線程作為在主機hostC程式的協助者。也就是說，核函數運作在GPU上而剩下的C代碼運作在CPU上。

CUDA程式設計模型也同時假定了主機和裝置同時在DRAM上分開持有他們自己的記憶體空間， 分别稱作host memory和device memeory。是以，一個程式同構調用CUDA 運作時（runtime）來管理着對所有線程可見的全局、常量和紋理記憶體空間。這些包括對裝置記憶體的配置設定和釋放以及主機和裝置之間的資料傳輸。

統一記憶體通過managed memory來連接配接主機和裝置兩個記憶體空間。managed memory可以通過一個使用共同的尋址空間來被系統中的所有GPUs和CPUs所通路。這就使得裝置記憶體使用可以“過載”且可以通過取消資料的在裝置端和主機端進行鏡像以極大地簡化應用的移植難度。

#計算力

裝置的計算力compute capability一般是由一個版本号表示，有時候也稱作SM version。 這個版本号揭示了GPU硬體所支援的特性和決定了應用程式在目前GPU運作時哪些特性或者指令是可用的。

計算力的版本号通過一個主版本X和一個小版本Y構成， 即 X.Y。擁有相同主版本的裝置具有通用的核心架構， 其中7表示Volta架構， 6表示Pascal架構， 5表示Maxwell架構， 3表示Kepler架構， 2表示Fermi架構， 1表示Tesla架構。小版本表示在主架構上進行了一些小提升，或者是一些新特性。