GPU(CUDA)学习日记(九)------ CUDA存储器模型
GPU(CUDA)學(xué)習(xí)日記(九)------ CUDA存儲(chǔ)器模型
標(biāo)簽: cuda存儲(chǔ)bindingcache編程api 2012-09-27 10:53 1677人閱讀 評(píng)論(1) 收藏 舉報(bào) 分類: GPU(16) CUDA(16) 高性能計(jì)算(7) 動(dòng)態(tài)加載(12) 并行運(yùn)算(7)CUDA存儲(chǔ)器模型:
GPU片內(nèi):register,shared memory;
板載顯存:local memory,constant memory, texture memory, texture memory,global memory;
host 內(nèi)存: host memory, pinned memory.
?
register: 訪問延遲極低;
????????????? 基本單元:register file (32bit/each)
????????????? 計(jì)算能力1.0/1.1版本硬件:8192/SM;
????????????? 計(jì)算能力1.2/1.3版本硬件: 16384/SM;
????????????? 每個(gè)線程占有的register有限,編程時(shí)不要為其分配過多私有變量;
local memory:寄存器被使用完畢,數(shù)據(jù)將被存儲(chǔ)在局部存儲(chǔ)器中;
? ? ? ? ? ? ? 大型結(jié)構(gòu)體或者數(shù)組;
? ? ? ? ? ? ? 無法確定大小的數(shù)組;
? ? ? ? ? ? ? 線程的輸入和中間變量;
? ? ? ? ? ? ? 定義線程私有數(shù)組的同時(shí)進(jìn)行初始化的數(shù)組被分配在寄存器中;
shared memory:訪問速度與寄存器相似;
? ? ? ? ? ? ? ?實(shí)現(xiàn)線程間通信的延遲最小;
? ? ? ? ? ? ? ?保存公用的計(jì)數(shù)器或者block的公用結(jié)果;
? ? ? ? ? ? ? ?硬件1.0~1.3中,16KByte/SM,被組織為16個(gè)bank;
? ? ? ? ? ? ? ?聲明關(guān)鍵字 _shared_? int sdata_static[16];
global memory:存在于顯存中,也稱為線性內(nèi)存(顯存可以被定義為線性存儲(chǔ)器或者CUDA數(shù)組);
? ? ? ? ? ? ? ?cudaMalloc()函數(shù)分配,cudaFree()函數(shù)釋放,cudaMemcpy()進(jìn)行主機(jī)端與設(shè)備端的數(shù)據(jù)傳輸;
? ? ? ? ? ? ? ?初始化共享存儲(chǔ)器需要調(diào)用cudaMemset();
? ? ? ? ? ? ? ?二維三維數(shù)組:cudaMallocPitch()和cudaMalloc3D()分配線性存儲(chǔ)空間,可以確保分配滿足對(duì)齊要求;
? ? ? ? ? ? ? cudaMemcpy2D(),cudaMemcpy3D()與設(shè)備端存儲(chǔ)器進(jìn)行拷貝;
?
host內(nèi)存:分為pageable memory 和 pinned memory
pageable memory: 通過操作系統(tǒng)API(malloc(),new())分配的存儲(chǔ)器空間;、
pinned memory:始終存在于物理內(nèi)存中,不會(huì)被分配到低速的虛擬內(nèi)存中,能夠通過DMA加速與設(shè)備端進(jìn)行通信;
? ? ? ? ? ? ? ? cudaHostAlloc(), cudaFreeHost()來分配和釋放pinned memory;
? ? ? ? ? ? ? ? 使用pinned memory優(yōu)點(diǎn):主機(jī)端-設(shè)備端的數(shù)據(jù)傳輸帶寬高;
? ? ? ? ? ? ? ? 某些設(shè)備上可以通過zero-copy功能映射到設(shè)備地址空間,從GPU直接訪問,省掉主存與顯存間進(jìn)行數(shù)據(jù)拷貝的工作;
? ? ? ? ? ? ? ?pinned memory 不可以分配過多:導(dǎo)致操作系統(tǒng)用于分頁的物理內(nèi)存變, 導(dǎo)致系統(tǒng)整體性能下降;通常由哪個(gè)cpu線程分配,就只有這個(gè)線程才有訪問權(quán)限;
? ? ? ? ? ? ? ? cuda2.3版本中,pinned memory功能擴(kuò)充:
? ? ? ? ? ? ? ?portable memory:讓控制不同GPU的主機(jī)端線程操作同一塊portable memory,實(shí)現(xiàn)cpu線程間通信;使用cudaHostAlloc()分配頁鎖定內(nèi)存時(shí),加上cudaHostAllocPortable標(biāo)志;
? ? ? ? ? ? ? ? write-combined Memory:提高從cpu向GPU單向傳輸數(shù)據(jù)的速度;不使用cpu的L1,L2 cache對(duì)一塊pinned memory中的數(shù)據(jù)進(jìn)行緩沖,將cache資源留給其他程序使用;在pci-e總線傳輸期間不會(huì)被來自cpu的監(jiān)視打斷;在調(diào)用cudaHostAlloc()時(shí)加上cudaHostAllocWriteCombined標(biāo)志;cpu從這種存儲(chǔ)器上讀取的速度很低;
? ? ? ? ? ? ? ? mapped memory:兩個(gè)地址:主機(jī)端地址(內(nèi)存地址),設(shè)備端地址(顯存地址)。??可以在kernnel程序中直接訪問mapped memory中的數(shù)據(jù),不必在內(nèi)存和顯存之間進(jìn)行數(shù)據(jù)拷貝,即zero-copy功能;在主機(jī)端可以由cudaHostAlloc()函數(shù)獲得,在設(shè)備端指針可以通過cudaHostGetDevicePointer()獲得;通過cudaGetDeviceProperties()函數(shù)返回的canMapHostMemory屬性知道設(shè)備是否支持mapped memory;在調(diào)用cudaHostAlloc()時(shí)加上cudaHostMapped標(biāo)志,將pinned memory映射到設(shè)備地址空間;必須使用同步來保證cpu和GPu對(duì)同一塊存儲(chǔ)器操作的順序一致性;顯存中的一部分可以既是portable memory又是mapped memory;在執(zhí)行CUDA操作前,先調(diào)用cudaSetDeviceFlags()(加cudaDeviceMapHost標(biāo)志)進(jìn)行頁鎖定內(nèi)存映射。
?
constant memory:只讀地址空間;位于顯存,有緩存加速;64Kb;用于存儲(chǔ)需要頻繁訪問的只讀參數(shù)?;只讀;使用_constant_ 關(guān)鍵字,定義在所有函數(shù)之外;兩種常數(shù)存儲(chǔ)器的使用方法:直接在定義時(shí)初始化常數(shù)存儲(chǔ)器;定義一個(gè)constant數(shù)組,然后使用函數(shù)進(jìn)行賦值;
?
texture memory:只讀;不是一塊專門的存儲(chǔ)器,而是牽涉到顯存、兩級(jí)紋理緩存、紋理拾取單元的紋理流水線;數(shù)據(jù)常以一維、二維或者三維數(shù)組的形式存儲(chǔ)在顯存中;緩存加速;可以聲明大小比常數(shù)存儲(chǔ)器大得多;適合實(shí)現(xiàn)圖像樹立和查找表;對(duì)大量數(shù)據(jù)的隨機(jī)訪問或非對(duì)齊訪問有良好的加速效果;在kernel中訪問紋理存儲(chǔ)器的操作成為紋理拾取(texture fetching);紋理拾取使用的坐標(biāo)與數(shù)據(jù)在顯存中的位置可以不同,通過紋理參照系約定二者的映射方式;將顯存中的數(shù)據(jù)與紋理參照系關(guān)聯(lián)的操作,稱為將數(shù)據(jù)與紋理綁定(texture binding);顯存中可以綁定到紋理的數(shù)據(jù)有:普通線性存儲(chǔ)器和cuda數(shù)組;存在緩存機(jī)制;可以設(shè)定濾波模式,尋址模式等;
本文的原文出處為:http://blog.csdn.net/ouczoe/article/details/5125621 幫助很大!謝謝原創(chuàng)作者!
與50位技術(shù)專家面對(duì)面20年技術(shù)見證,附贈(zèng)技術(shù)全景圖總結(jié)
以上是生活随笔為你收集整理的GPU(CUDA)学习日记(九)------ CUDA存储器模型的全部?jī)?nèi)容,希望文章能夠幫你解決所遇到的問題。
- 上一篇: MEGA | 多序列比对及系统发育树的构
- 下一篇: 高颜值在线绘图平台ImageGP系列教程