CUDA並行存儲模型

CUDA將CPU做爲主機（Host），GPU做爲設備（Device）。一個系統中能夠有一個主機和多個設備。CPU負責邏輯性強的事務處理和串行計算，GPU專一於執行高度線程化的並行處理任務。它們擁有相互獨立的存儲器（主機端的內存和顯卡端的顯存）。 

 

運行在GPU上的函數稱爲kernel（內核函數）。一個完整的CUDA程序是由一些列的kernel函數和主機端的串行處理步驟共同完成的。CPU串行代碼的工做包括在kernel啓動前進行的數據準備、設備初始化以及在kernel之間進行一些串行化計算。 

 

kernel函數以函數類型限定符_global_定義，而且只能在主機端代碼中調用。調用時須要制定kernel的Grid中的block數目以及每一個block中thread的數目。每一個線程有本身的blockID和threadID，可用來區分其它的線程，這兩個內建變量是隻讀的，由專用寄存器提供，只能有kernel函數調用。 

 

CUDA將計算任務映射爲大量的可並行執行的線程，並由硬件動態調度和執行這些線程。kernel是以線程網格Grid的形式組織，每一個Grid又若干個線程塊block組成，每一個block又由若干線程thread組成。本質上，kernel是以block爲單位執行的，grid只是用來表示一些列並行block的集合。Block之間是不能彼此通訊的。目前一個kernel只支持一個grid，在多指令多數據（MIMD）構架中會存在多個grid。

 

爲方便編程，CUDA使用了dim3類型的內建變量threadIdx和threadIdx。 

對於一維的block，線程的編號是threadIdx.x。 

對於二維（Dx，Dy）的block，線程的編號是threadIdx.x+threadIdx.y * Dx。 

對於三維（Dx，Dy，Dz）的block，線程的編號是threadIdx.x+threadIdx.y * Dx+hreadIdx.z* Dx * Dy。 

 

GPU的計算核心是流多處理器（SM），每一個SM包含8個標量流處理器（SP）以及其它的運算單元。Kernel是以block爲單位執行的，同一個block的線程須要共享數據，所以它們共享同一個SM。一個block必須分配到一個SM，可是能夠一個SM中同一個時刻有多個活動塊（active block）執行等待，即同一個SM能夠有多個block上下文。實際運行中，block會被分爲更小的線程束（wrap），線程束的大小由硬件的計算能力決定，Tesla的構架中一個wrap由32個線程組成。 

CUDA採用了單指令多線程執行模型。這個模型是對單指令多數據的改進。CUDA中執行寬度能夠在1——512個線程之間變化，可是在單指令多數據中執行寬度必須是一個wrap（32）。

每個線程擁有本身的私有存儲器，每個線程塊擁有一塊共享存儲器(Shared memory)；最後，grid中全部的線程均可以訪問同一塊全局存儲器(global memory)。除此以外，還有兩種能夠被全部線程訪問的只讀存儲器：常數存儲器（constant memory）和紋理存儲器(Texture memory),它們分別爲不一樣的應用進行了優化。全局存儲器、常數存儲器和紋理存儲器中的值在一個內核函數執行完成後將被繼續保持，能夠被同一程序中其也內核函數調用。

存儲器	位置	擁有緩存	訪問權限	變量生存週期
register	GPU片內	N/A	Device可讀/寫	與thread相同
Local memory	板載顯存	無	Device可讀/寫	與thread相同
Shared memory	GPU片內	N/A	Device可讀/寫	與block相同
Constant memory	板載顯存	有	Device可讀，host要讀寫	可在程序中保持
Texture memory	板載顯存	有	Device可讀，host要讀寫	可在程序中保持
Global memory	板載顯存	無	Device可讀/寫, host可讀/寫	可在程序中保持
Host memory	Host內存	無	host可讀/寫	可在程序中保持
Pinned memory	Host內存	無	host可讀/寫	可在程序中保持

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

 

CUDA存儲器模型：

GPU片內：register,shared memory;

板載顯存：local memory,constant 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數組；存在緩存機制；能夠設定濾波模式，尋址模式等；