簡述CUDA執行緒及求CUDA中執行緒索引

CUDA 簡介

2006 年，NVIDIA公司釋出了CUDA（Compute Unified Device Architecture)，是一種新的操作GPU計算的硬體和軟體架構，是建立在NVIDIA的GPUs上的一個通用平行計算平臺和程式設計模型，它提供了GPU編理的簡易介面，基於CUDA程式設計可以構建基於GPU計算的應用程式，利用GPUs的平行計算引擎來更加高效地解決比較複雜的計算難驗。它將GPU視作一個資料並行裝置，而且無需把這些計算對映影像API。作業系統的多工機制可以同時管理CUDA訪問GPU和圖形程式的執行庫，其計算特性支援利用CUDA直觀編寫GPU核心程式。

CUDA 在軟體方面組成有：一個CUDA庫，一個應用程式程式設計介面(API)及其執行庫(Runtime)，兩個較高階別的通用數學庫，即CUFFT和CUBLAS，CUDA改講了DRAM的讀寫靈活性，使得GPU與CPU的機制相吻合。另一方面，CUDA提供了片上(on-chip)共享記憶體，使得執行緒之間可以共享資料。應用程式可以利用共享記憶體來減少DRAM的資料傳送，更少的依賴DRAM的記憶體頻寬。

 

CUDA 執行緒模型

CUDA 的架構中引入了主機端（host）和裝置（device）的概念。CUDA程式中既包含host程式，又包含device程式。同時，host與device之間可以進行通訊，這樣它們之間可以進行資料複製。

主機(Host)：將CPU及系統的記憶體（記憶體條）稱為主機。

裝置(Device)：將GPU及GPU本身的顯示記憶體稱為裝置。

典型的CUDA程式的執行流程如下：

1. 分配host記憶體，並進行資料初始化；

2. 分配device記憶體：  cudaMalloc(void **devPtr, size_t count);

3. 並從host將資料複製到device上：cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyHostToDevice)；

4. 呼叫CUDA的核函式在device上完成指定的運算：<<<M, T>>>；

   M 表示kernel launches with a gird of M thread blocks。T表示每個thread block具有T parallel threads。

5. 將device上的運算結果複製到host上：cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyDeviceToHost)；

6. 釋放device和host上分配的記憶體：cudaFree（），free（）。 

 

CUDA 執行緒層次結構

 

1. 核 kernel

CUDA 執行流程中最重要的一個過程是呼叫CUDA的核函式來執行平行計算，kernel是CUDA中一個重要的概念。在CUDA程式構架中，主機端程式碼部分在CPU上執行，是普通的C程式碼；當遇到資料並行處理的部分，CUDA 就會將程式編譯成GPU能執行的程式，並傳送到GPU，這個程式在CUDA裡稱做核（kernel）。裝置端程式碼部分在GPU上執行，此程式碼部分在kernel上編寫（.cu檔案）。kernel用global符號宣告，在呼叫時需要用<<<grid, block>>>來指定kernel要執行及結構。

CUDA 是透過函式型別限定詞區別在host和device上的函式，主要的三個函式型別限定詞如下：

       i.          global ：在device上執行，從host中呼叫（一些特定的GPU也可以從device上呼叫），返回型別必須是void，不支援可變引數引數，不能成為類成員函式。注意用global定義的kernel是非同步的，這意味著host不會等待kernel執行完就執行下一步。

     ii.          device ：在device上執行，單僅可以從device中呼叫，不可以和global同時用。

    iii.          host ：在host上執行，僅可以從host上呼叫，一般省略不寫，不可以和global同時用，但可和device同時使用，此時函式會在device和host都編譯。

 

2. 網格 grid

kernel 在device上執行時，實際上是啟動很多執行緒，一個kernel所啟動的所有執行緒稱為一個網格（grid），同一個網格上的執行緒共享相同的全域性記憶體空間。grid是執行緒結構的第一層次。

 

3. 執行緒塊 block

網格又可以分為很多執行緒塊（block），一個block裡面包含很多執行緒。各block是並行執行的，block間無法通訊，也沒有執行順序。block的數量限制為不超過65535（硬體限制）。block是執行緒結構的第二層次。

grid 和block都是定義為dim3型別的變數，dim3可以看成是包含三個無符號整數（x，y，z）成員的結構體變數，在定義時，預設值初始化為1。grid和block可以靈活地定義為1-dim，2-dim以及3-dim結構。

CUDA 中，每一個執行緒都要執行核函式，每一個執行緒需要kernel的兩個內建座標變數（blockIdx，threadIdx）來唯一標識，其中blockIdx指明執行緒所在grid中的位置，threaIdx指明執行緒所在block中的位置。它們都是dim3型別變數。

 

4. 執行緒 thread

一個CUDA的並行程式會被以許多個threads來執行。數個threads會被群組成一個block，同一個block中的threads可以同步，也可以透過shared memory通訊。

 

求CUDA中執行緒索引

可以把網格和執行緒塊都看作一個三維的矩陣。這裡假設網格是一個345的三維矩陣， 執行緒塊是一個456的三維矩陣。

gridDim

gridDim.x 、gridDim.y、gridDim.z分別表示網格各個維度的大小，所以有

gridDim.x=3

gridDim.y=4

gridDim.z=5

 

blockDim

blockDim.x 、blockDim.y、blockDim.z分別表示執行緒塊中各個維度的大小，所以有

blockDim.x=4

blockDim.y=5

blockDim.z=6

 

blockIdx

blockIdx.x 、blockIdx.y、blockIdx.z分別表示當前執行緒塊所處的執行緒格的座標位置

 

threadIdx

threadIdx.x 、threadIdx.y、threadIdx.z分別表示當前執行緒所處的執行緒塊的座標位置

 

網格里面總的執行緒個數N即可透過下面的公式算出

N = gridDim.x*gridDim.y*gridDim.z*blockDim.x*blockDim.y*blockDim.z

 

舉例：

將所有的執行緒排成一個序列，序列號為0 , 1 , 2 , … , N ，如何找到當前的序列號？

1. 先找到當前執行緒位於網格中的哪一個執行緒塊blockId

   blockId = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;

2. 找到當前執行緒位於執行緒塊中的哪一個執行緒threadId

   threadId = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;

3. 計算一個執行緒塊中一共有多少個執行緒M

   M = blockDim.x*blockDim.y*blockDim.z

4. 求得當前的執行緒序列號idx

   idx = threadId + M*blockId;

本文轉載自公眾號：BFT智慧機器人研究

https://mp.weixin.qq.com/s/_3RSg_b0Q7HcqPdibCE6Qg