GPU编程之GUDA(四)——基础知识补充
?
主機函數:在CPU上調用,CPU上執行的函數
全局函數:在CPU上調用,GPU上執行的函數
設備函數:在GPU上調用,GPU上執行的函數——它的線程配置由調用關系中最近的全局函數決定
?
主機函數在聲明時可以帶有限定符 _host_
全局函數在聲明時必須帶有限定符 _global_
設備函數在聲明時必須帶有限定符 _device_
?
在調用全局函數時,除了函數名和形參表外,還有一個用三個小于號“<”和三個大于號“>”包含的部分,這一部分用來指定在并行計算式使用的線程組數和每個線程組白虎呢的線程數。
my_first_kernel<<<nBlocks,nThreads>>>(/*全局函數參數表*/)
一個grid包含完成一次函數調用的所有block,一個block由若干個線程組成
block和grid的尺寸都可以用三元向量來表示,這表明block和grid都是三位數組
?
grid的數組元素由block組成,block的數組元素由線程組成
?
表示線程數量時,可以用CUDA自帶的dim3變量來聲明三元向量,如果是二元向量,即和第三元置為一是等效的——可缺省,如果只是一維的,可以用int代替
dim3 gridsize(3,2,1); dim3 blocksize(2,2,2); my_first_kernel<<<gridsize,blocksize>>>();dim3 gridsize(3,2);?
內建向量是從基本體數據類型擴展而成的結構體,訪問內建向量的元素可以用索引x,y.z.w
?
=======================================================================================================
運用GUDA運算時,與經典GPGPU方法一樣,需要一個輸入緩存來準備原始數據和一個輸出緩存來接收GPU計算的結果
因此,需要為程序分配兩個數組:一個設備上的數組pfDevice用來存放GPU生成的數據,一個在主機上的數組pfHost,用來接收從GU取回的數據
這兩個數組應該擁有相同的大小和類型
=======================================================================================================
?
線程索引,是線程在每個block里的索引,由于block的尺寸是三維的,因此線程索引也是一個三元向量:threadIdx
一個grid中的block索引:block.x? block.y? block.z
一個block中的thread索引:thread.x? thread.y? thread.z?
一次函數調用只有一個grid,因此不存在grid索引
?
block尺寸保存在向量blockDim中,grid尺寸保存在gridDim中
線程號:每個線程唯一的標識符
當grid和block是一維的時候,線程號是這樣直觀定義的
int tid=threadIdx.x+blockDim.x*blockIdx.x;?
在實際運用中,如果沒有某些需要,一般都沒有必要使用都為的block,因為一維的block下線程比較容易管理
什么是某些需要?
這樣的“某些需要”可以和硬件有關(1、某些存儲器為了二維元素而特別優化,2、紋理存儲器),也可以是算法的需要,有些算法在使用二維block時較為直觀,有助于用戶實現算法
=======================================================================================================
內核:人們習慣將全局函數這樣在GPU上執行的函數稱為內核,簡稱核。它是函數執行時的基本單位
?
一個grid對應一個GPU或者多個多處理器
一個block中的所有線程在一個多處理器上并發執行
對于同一個block中的線程來說:一、同一個block中的線程可以同步
?????????????????????????????????????????????????? ?二、他們可以共同訪問多處理器里的共享存儲器
??????????????????????????????????????????????????? 三:他們不能被拆分到多個多處理器上執行
對處理器將每個縣城與一個標量處理器核心對應起來
?
warp:32個線程
如果block中的線程數為32的整數倍,會給調度帶來方便。
?
CUDA存儲器分為三層:1、私有本地存儲器:可被單個線程訪問
???????????????????????????????????????2、共享存儲器:可被block中所有線程訪問
???????????????????????????????????????3、全局存儲器、本地存儲器:可被一個應用程序中所有線程訪問
除此之外還有程序中所有線程都可以訪問的只讀存儲器:常量存儲器和紋理存儲器
?
自動變量:在全局函數和設備函數中不是用限定符說明的變量——系統決定他們存放在哪里——————一般放在寄存器中
寄存器:獨立分布于貝格標量處理器上,為每個線程提供私有的存儲空間
??????????? 特點:容量小,訪問速度快。
?
共享存儲器:位于么一個多處理器內——兩種方式:靜態分配方式,動態分配方式(動態分配時,存儲器尺寸可以是變量)
??????????????????? 在調用全局函數時不需要在執行配置表中給出共享存儲器的尺寸,直接在內核中用限定符_shared_即可
?
對共享存儲器訪問是以半個warp為單位的
bank沖突:當半warp中的多個線程方位的數組元素位于同一個bank時,他們的讀寫操作就不能同時進行,也就是說,這幾個對同一bank的訪問會被串行化
?
無沖突訪問模式:1、分在不同的bank中,2、廣播
?
設備存儲器分為:全局存儲器、紋理存儲器、常量存儲器
設備存儲器有兩種分配方式:分配為線性存儲器、也可以被分配為CUDA數組
?
CUDA數組是為讀取紋理而特別優化的,因此一般將紋理存儲器分配為CUDA數組
全局存儲器最便捷的使用方式是線性存儲的方式
常量存儲器:容量小,只讀
?
設備存儲器不是安裝在多處理器內的存儲器,即片外存儲器,因此對設備存儲器的直接訪問需要耗費大量的時間
緩存、寄存器、共享存儲器訪問速度快
?
使用全局存儲器需要滿足接合的要求:不僅要求訪問連續的存儲單元,還對訪問的起始地址有一定的要求——數據的起始地址應為每個線程訪問數據大小的16倍的整數倍
一維、二維、四維向量可以結合成較大的數據傳輸單元,二三維向量的結核性能稍遜,只能作為三個一維標量來處理
?
并發:
數據傳輸與內核的并發
內核的并發
數據傳輸的并發
?
流:不同的流在執行時是相互獨立的,因此他們可能在時間軸上出現亂序的情況,即不是按照定義的順序執行,故用戶需要再設計時考慮到這一點可能帶來的隱患
?
參考文獻:
仇德元.《GPGPU編程技術——從GLSL、CDPU到OpenGL》[M].河北省三河市:機械工業出版社,2012年:323.
?
總結
以上是生活随笔為你收集整理的GPU编程之GUDA(四)——基础知识补充的全部內容,希望文章能夠幫你解決所遇到的問題。
- 上一篇: 租用服务器选择大带宽租用具备哪些优势
- 下一篇: 如何用word制作商品条码