?

主機函數(shù)：在CPU上調(diào)用，CPU上執(zhí)行的函數(shù)

全局函數(shù)：在CPU上調(diào)用，GPU上執(zhí)行的函數(shù)

設(shè)備函數(shù)：在GPU上調(diào)用，GPU上執(zhí)行的函數(shù)——它的線程配置由調(diào)用關(guān)系中最近的全局函數(shù)決定

?

主機函數(shù)在聲明時可以帶有限定符 _host_

全局函數(shù)在聲明時必須帶有限定符 _global_

設(shè)備函數(shù)在聲明時必須帶有限定符 _device_

?

在調(diào)用全局函數(shù)時，除了函數(shù)名和形參表外，還有一個用三個小于號“<”和三個大于號“>”包含的部分，這一部分用來指定在并行計算式使用的線程組數(shù)和每個線程組白虎呢的線程數(shù)。

my_first_kernel<<<nBlocks,nThreads>>>(/*全局函數(shù)參數(shù)表*/)

一個grid包含完成一次函數(shù)調(diào)用的所有block，一個block由若干個線程組成

block和grid的尺寸都可以用三元向量來表示，這表明block和grid都是三位數(shù)組

?

grid的數(shù)組元素由block組成，block的數(shù)組元素由線程組成

?

表示線程數(shù)量時，可以用CUDA自帶的dim3變量來聲明三元向量，如果是二元向量，即和第三元置為一是等效的——可缺省，如果只是一維的，可以用int代替

dim3 gridsize(3,2,1); dim3 blocksize(2,2,2); my_first_kernel<<<gridsize,blocksize>>>();dim3 gridsize(3,2);

?

內(nèi)建向量是從基本體數(shù)據(jù)類型擴展而成的結(jié)構(gòu)體，訪問內(nèi)建向量的元素可以用索引x,y.z.w

?

=======================================================================================================

運用GUDA運算時，與經(jīng)典GPGPU方法一樣，需要一個輸入緩存來準備原始數(shù)據(jù)和一個輸出緩存來接收GPU計算的結(jié)果

因此，需要為程序分配兩個數(shù)組：一個設(shè)備上的數(shù)組pfDevice用來存放GPU生成的數(shù)據(jù)，一個在主機上的數(shù)組pfHost，用來接收從GU取回的數(shù)據(jù)

這兩個數(shù)組應(yīng)該擁有相同的大小和類型

=======================================================================================================

?

線程索引，是線程在每個block里的索引，由于block的尺寸是三維的，因此線程索引也是一個三元向量：threadIdx

一個grid中的block索引：block.x? block.y? block.z

一個block中的thread索引：thread.x? thread.y? thread.z?

一次函數(shù)調(diào)用只有一個grid，因此不存在grid索引

?

block尺寸保存在向量blockDim中，grid尺寸保存在gridDim中

線程號：每個線程唯一的標識符

當grid和block是一維的時候，線程號是這樣直觀定義的

int tid=threadIdx.x+blockDim.x*blockIdx.x;

?

在實際運用中，如果沒有某些需要，一般都沒有必要使用都為的block，因為一維的block下線程比較容易管理

什么是某些需要？

這樣的“某些需要”可以和硬件有關(guān)（1、某些存儲器為了二維元素而特別優(yōu)化，2、紋理存儲器），也可以是算法的需要，有些算法在使用二維block時較為直觀，有助于用戶實現(xiàn)算法

=======================================================================================================

內(nèi)核：人們習(xí)慣將全局函數(shù)這樣在GPU上執(zhí)行的函數(shù)稱為內(nèi)核，簡稱核。它是函數(shù)執(zhí)行時的基本單位

?

一個grid對應(yīng)一個GPU或者多個多處理器

一個block中的所有線程在一個多處理器上并發(fā)執(zhí)行

對于同一個block中的線程來說：一、同一個block中的線程可以同步

?????????????????????????????????????????????????? ?二、他們可以共同訪問多處理器里的共享存儲器

??????????????????????????????????????????????????? 三：他們不能被拆分到多個多處理器上執(zhí)行

對處理器將每個縣城與一個標量處理器核心對應(yīng)起來

?

warp：32個線程

如果block中的線程數(shù)為32的整數(shù)倍，會給調(diào)度帶來方便。

?

CUDA存儲器分為三層：1、私有本地存儲器：可被單個線程訪問

???????????????????????????????????????2、共享存儲器：可被block中所有線程訪問

???????????????????????????????????????3、全局存儲器、本地存儲器：可被一個應(yīng)用程序中所有線程訪問

除此之外還有程序中所有線程都可以訪問的只讀存儲器：常量存儲器和紋理存儲器

?

自動變量：在全局函數(shù)和設(shè)備函數(shù)中不是用限定符說明的變量——系統(tǒng)決定他們存放在哪里——————一般放在寄存器中

寄存器：獨立分布于貝格標量處理器上，為每個線程提供私有的存儲空間

??????????? 特點：容量小，訪問速度快。

?

共享存儲器：位于么一個多處理器內(nèi)——兩種方式：靜態(tài)分配方式，動態(tài)分配方式（動態(tài)分配時，存儲器尺寸可以是變量）

??????????????????? 在調(diào)用全局函數(shù)時不需要在執(zhí)行配置表中給出共享存儲器的尺寸，直接在內(nèi)核中用限定符_shared_即可

?

對共享存儲器訪問是以半個warp為單位的

bank沖突：當半warp中的多個線程方位的數(shù)組元素位于同一個bank時，他們的讀寫操作就不能同時進行，也就是說，這幾個對同一bank的訪問會被串行化

?

無沖突訪問模式：1、分在不同的bank中，2、廣播

?

設(shè)備存儲器分為：全局存儲器、紋理存儲器、常量存儲器

設(shè)備存儲器有兩種分配方式：分配為線性存儲器、也可以被分配為CUDA數(shù)組

?

CUDA數(shù)組是為讀取紋理而特別優(yōu)化的，因此一般將紋理存儲器分配為CUDA數(shù)組

全局存儲器最便捷的使用方式是線性存儲的方式

常量存儲器：容量小，只讀

?

設(shè)備存儲器不是安裝在多處理器內(nèi)的存儲器，即片外存儲器，因此對設(shè)備存儲器的直接訪問需要耗費大量的時間

緩存、寄存器、共享存儲器訪問速度快

?

使用全局存儲器需要滿足接合的要求：不僅要求訪問連續(xù)的存儲單元，還對訪問的起始地址有一定的要求——數(shù)據(jù)的起始地址應(yīng)為每個線程訪問數(shù)據(jù)大小的16倍的整數(shù)倍

一維、二維、四維向量可以結(jié)合成較大的數(shù)據(jù)傳輸單元，二三維向量的結(jié)核性能稍遜，只能作為三個一維標量來處理

?

并發(fā)：

數(shù)據(jù)傳輸與內(nèi)核的并發(fā)

內(nèi)核的并發(fā)

數(shù)據(jù)傳輸?shù)牟l(fā)

?

流：不同的流在執(zhí)行時是相互獨立的，因此他們可能在時間軸上出現(xiàn)亂序的情況，即不是按照定義的順序執(zhí)行，故用戶需要再設(shè)計時考慮到這一點可能帶來的隱患

?

參考文獻：

仇德元．《GPGPU編程技術(shù)——從GLSL、CDPU到OpenGL》[M]．河北省三河市：機械工業(yè)出版社，2012年：323．

?

總結(jié)

以上是生活随笔為你收集整理的GPU编程之GUDA（四）——基础知识补充的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。