CUDA存儲器模型:
GPU片内:register,shared memory;
闆載顯存:local memory,constant memory, texture memory, texture memory,global memory;
host 記憶體: host memory, pinned memory.
register: 通路延遲極低;
基本單元:register file (32bit/each)
計算能力1.0/1.1版本硬體:8192/SM;
計算能力1.2/1.3版本硬體: 16384/SM;
每個線程占有的register有限,程式設計時不要為其配置設定過多私有變量;
local memory:寄存器被使用完畢,資料将被存儲在局部存儲器中;
大型結構體或者數組;
無法确定大小的數組;
線程的輸入和中間變量;
定義線程私有數組的同時進行初始化的數組被配置設定在寄存器中;
shared memory:通路速度與寄存器相似;
實作線程間通信的延遲最小;
儲存公用的計數器或者block的公用結果;
硬體1.0~1.3中,16KByte/SM,被組織為16個bank;
聲明關鍵字 _shared_ int sdata_static[16];
global memory:存在于顯存中,也稱為線性記憶體(顯存可以被定義為線性存儲器或者CUDA數組);
cudaMalloc()函數配置設定,cudaFree()函數釋放,cudaMemcpy()進行主機端與裝置端的資料傳輸;
初始化共享存儲器需要調用cudaMemset();
二維三維數組:cudaMallocPitch()和cudaMalloc3D()配置設定線性存儲空間,可以確定配置設定滿足對齊要求;
cudaMemcpy2D(),cudaMemcpy3D()與裝置端存儲器進行拷貝;
host記憶體:分為pageable memory 和 pinned memory
pageable memory: 通過作業系統API(malloc(),new())配置設定的存儲器空間;、
pinned memory:始終存在于實體記憶體中,不會被配置設定到低速的虛拟記憶體中,能夠通過DMA加速與裝置端進行通信;
cudaHostAlloc(), cudaFreeHost()來配置設定和釋放pinned memory;
使用pinned memory優點:主機端-裝置端的資料傳輸帶寬高;
某些裝置上可以通過zero-copy功能映射到裝置位址空間,從GPU直接通路,省掉主存與顯存間進行資料拷貝的工作;
pinned memory 不可以配置設定過多:導緻作業系統用于分頁的實體記憶體變, 導緻系統整體性能下降;通常由哪個cpu線程配置設定,就隻有這個線程才有通路權限;
cuda2.3版本中,pinned memory功能擴充:
portable memory:讓控制不同GPU的主機端線程操作同一塊portable memory,實作cpu線程間通信;使用cudaHostAlloc()配置設定頁鎖定記憶體時,加上cudaHostAllocPortable标志;
write-combined Memory:提高從cpu向GPU單向傳輸資料的速度;不使用cpu的L1,L2 cache對一塊pinned memory中的資料進行緩沖,将cache資源留給其他程式使用;在pci-e總線傳輸期間不會被來自cpu的監視打斷;在調用cudaHostAlloc()時加上cudaHostAllocWriteCombined标志;cpu從這種存儲器上讀取的速度很低;
mapped memory:兩個位址:主機端位址(記憶體位址),裝置端位址(顯存位址)。 可以在kernnel程式中直接通路mapped memory中的資料,不必在記憶體和顯存之間進行資料拷貝,即zero-copy功能;在主機端可以由cudaHostAlloc()函數獲得,在裝置端指針可以通過cudaHostGetDevicePointer()獲得;通過cudaGetDeviceProperties()函數傳回的canMapHostMemory屬性知道裝置是否支援mapped memory;在調用cudaHostAlloc()時加上cudaHostMapped标志,将pinned memory映射到裝置位址空間;必須使用同步來保證cpu和GPu對同一塊存儲器操作的順序一緻性;顯存中的一部分可以既是portable memory又是mapped memory;在執行CUDA操作前,先調用cudaSetDeviceFlags()(加cudaDeviceMapHost标志)進行頁鎖定記憶體映射。
constant memory:隻讀位址空間;位于顯存,有緩存加速;64Kb;用于存儲需要頻繁通路的隻讀參數 ;隻讀;使用_constant_ 關鍵字,定義在所有函數之外;兩種常數存儲器的使用方法:直接在定義時初始化常數存儲器;定義一個constant數組,然後使用函數進行指派;
texture memory:隻讀;不是一塊專門的存儲器,而是牽涉到顯存、兩級紋理緩存、紋理拾取單元的紋理流水線;資料常以一維、二維或者三維數組的形式存儲在顯存中;緩存加速;可以聲明大小比常數存儲器大得多;适合實作圖像樹立和查找表;對大量資料的随機通路或非對齊通路有良好的加速效果;在kernel中通路紋理存儲器的操作成為紋理拾取(texture fetching);紋理拾取使用的坐标與資料在顯存中的位置可以不同,通過紋理參照系約定二者的映射方式;将顯存中的資料與紋理參照系關聯的操作,稱為将資料與紋理綁定(texture binding);顯存中可以綁定到紋理的資料有:普通線性存儲器和cuda數組;存在緩存機制;可以設定濾波模式,尋址模式等;
本文的原文出處為:http://blog.csdn.net/ouczoe/article/details/5125621