CUDA簡(jiǎn)介

2006年，NVIDIA公司發(fā)布了CUDA（Compute Unified Device Architecture)，是一種新的操作GPU計(jì)算的硬件和軟件架構(gòu)，是建立在NVIDIA的GPUs上的一個(gè)通用并行計(jì)算平臺(tái)和編程模型，它提供了GPU編理的簡(jiǎn)易接口，基于CUDA編程可以構(gòu)建基于GPU計(jì)算的應(yīng)用程序，利用GPUs的并行計(jì)算引擎來更加高效地解決比較復(fù)雜的計(jì)算難驗(yàn)。它將GPU視作一個(gè)數(shù)據(jù)并行設(shè)備，而且無需把這些計(jì)算映射圖像API。操作系統(tǒng)的多任務(wù)機(jī)制可以同時(shí)管理CUDA訪問GPU和圖形程序的運(yùn)行庫，其計(jì)算特性支持利用CUDA直觀編寫GPU核心程序。

CUDA在軟件方面組成有：一個(gè)CUDA庫，一個(gè)應(yīng)用程序編程接口(API)及其運(yùn)行庫(Runtime)，兩個(gè)較高級(jí)別的通用數(shù)學(xué)庫，即CUFFT和CUBLAS，CUDA改講了DRAM的讀寫靈活性，使得GPU與CPU的機(jī)制相吻合。另一方面，CUDA提供了片上(on-chip)共享內(nèi)存，使得線程之間可以共享數(shù)據(jù)。應(yīng)用程序可以利用共享內(nèi)存來減少DRAM的數(shù)據(jù)傳送，更少的依賴DRAM的內(nèi)存帶寬。

?

CUDA線程模型

CUDA的架構(gòu)中引入了主機(jī)端（host）和設(shè)備（device）的概念。CUDA程序中既包含host程序，又包含device程序。同時(shí)，host與device之間可以進(jìn)行通信，這樣它們之間可以進(jìn)行數(shù)據(jù)拷貝。

主機(jī)(Host)：將CPU及系統(tǒng)的內(nèi)存（內(nèi)存條）稱為主機(jī)。

設(shè)備(Device)：將GPU及GPU本身的顯示內(nèi)存稱為設(shè)備。

典型的CUDA程序的執(zhí)行流程如下：

1. 分配host內(nèi)存，并進(jìn)行數(shù)據(jù)初始化；

2. 分配device內(nèi)存：? cudaMalloc(void **devPtr, size_t count);

3. 并從host將數(shù)據(jù)拷貝到device上：cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyHostToDevice)；

4. 調(diào)用CUDA的核函數(shù)在device上完成指定的運(yùn)算：<<<M, T>>>；

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

5. 將device上的運(yùn)算結(jié)果拷貝到host上：cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyDeviceToHost)；

6. 釋放device和host上分配的內(nèi)存：cudaFree（），free（）。?

?

CUDA線程層次結(jié)構(gòu)

1. 核 kernel

CUDA執(zhí)行流程中最重要的一個(gè)過程是調(diào)用CUDA的核函數(shù)來執(zhí)行并行計(jì)算，kernel是CUDA中一個(gè)重要的概念。在CUDA程序構(gòu)架中，主機(jī)端代碼部分在CPU上執(zhí)行，是普通的C代碼；當(dāng)遇到數(shù)據(jù)并行處理的部分，CUDA 就會(huì)將程序編譯成GPU能執(zhí)行的程序，并傳送到GPU，這個(gè)程序在CUDA里稱做核（kernel）。設(shè)備端代碼部分在GPU上執(zhí)行，此代碼部分在kernel上編寫（.cu文件）。kernel用global符號(hào)聲明，在調(diào)用時(shí)需要用<<<grid, block>>>來指定kernel要執(zhí)行及結(jié)構(gòu)。

CUDA是通過函數(shù)類型限定詞區(qū)別在host和device上的函數(shù)，主要的三個(gè)函數(shù)類型限定詞如下：

?????i.?global：在device上執(zhí)行，從host中調(diào)用（一些特定的GPU也可以從device上調(diào)用），返回類型必須是void，不支持可變參數(shù)參數(shù)，不能成為類成員函數(shù)。注意用global定義的kernel是異步的，這意味著host不會(huì)等待kernel執(zhí)行完就執(zhí)行下一步。

???? ii.?device：在device上執(zhí)行，單僅可以從device中調(diào)用，不可以和global同時(shí)用。

?????iii.?host：在host上執(zhí)行，僅可以從host上調(diào)用，一般省略不寫，不可以和global同時(shí)用，但可和device同時(shí)使用，此時(shí)函數(shù)會(huì)在device和host都編譯。

?

2. 網(wǎng)格 grid

kernel在device上執(zhí)行時(shí)，實(shí)際上是啟動(dòng)很多線程，一個(gè)kernel所啟動(dòng)的所有線程稱為一個(gè)網(wǎng)格（grid），同一個(gè)網(wǎng)格上的線程共享相同的全局內(nèi)存空間。grid是線程結(jié)構(gòu)的第一層次。

?

3. 線程塊 block

網(wǎng)格又可以分為很多線程塊（block），一個(gè)block里面包含很多線程。各block是并行執(zhí)行的，block間無法通信，也沒有執(zhí)行順序。block的數(shù)量限制為不超過65535（硬件限制）。block是線程結(jié)構(gòu)的第二層次。

grid和block都是定義為dim3類型的變量，dim3可以看成是包含三個(gè)無符號(hào)整數(shù)（x，y，z）成員的結(jié)構(gòu)體變量，在定義時(shí)，缺省值初始化為1。grid和block可以靈活地定義為1-dim，2-dim以及3-dim結(jié)構(gòu)。

CUDA中，每一個(gè)線程都要執(zhí)行核函數(shù)，每一個(gè)線程需要kernel的兩個(gè)內(nèi)置坐標(biāo)變量（blockIdx，threadIdx）來唯一標(biāo)識(shí)，其中blockIdx指明線程所在grid中的位置，threaIdx指明線程所在block中的位置。它們都是dim3類型變量。

?

4. 線程 thread

一個(gè)CUDA的并行程序會(huì)被以許多個(gè)threads來執(zhí)行。數(shù)個(gè)threads會(huì)被群組成一個(gè)block，同一個(gè)block中的threads可以同步，也可以通過shared memory通信。

?

求CUDA中線程索引

可以把網(wǎng)格和線程塊都看作一個(gè)三維的矩陣。這里假設(shè)網(wǎng)格是一個(gè)345的三維矩陣， 線程塊是一個(gè)456的三維矩陣。

gridDim

gridDim.x、gridDim.y、gridDim.z分別表示網(wǎng)格各個(gè)維度的大小，所以有

gridDim.x=3

gridDim.y=4

gridDim.z=5

?

blockDim

blockDim.x、blockDim.y、blockDim.z分別表示線程塊中各個(gè)維度的大小，所以有

blockDim.x=4

blockDim.y=5

blockDim.z=6

?

blockIdx

blockIdx.x、blockIdx.y、blockIdx.z分別表示當(dāng)前線程塊所處的線程格的坐標(biāo)位置

?

threadIdx

threadIdx.x、threadIdx.y、threadIdx.z分別表示當(dāng)前線程所處的線程塊的坐標(biāo)位置

?

網(wǎng)格里面總的線程個(gè)數(shù)N即可通過下面的公式算出

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

?

舉例：

將所有的線程排成一個(gè)序列，序列號(hào)為0 , 1 , 2 , … , N ，如何找到當(dāng)前的序列號(hào)？

1. 先找到當(dāng)前線程位于網(wǎng)格中的哪一個(gè)線程塊blockId

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

2. 找到當(dāng)前線程位于線程塊中的哪一個(gè)線程threadId

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

3. 計(jì)算一個(gè)線程塊中一共有多少個(gè)線程M

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

4. 求得當(dāng)前的線程序列號(hào)idx

?? idx = threadId + M*blockId;

本文轉(zhuǎn)載自公眾號(hào)：BFT智能機(jī)器人研究

更多機(jī)器人信息等你來看~