CUDA儲存空間模型

來源:互聯網
上載者:User

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對同一Block Storage器操作的順序一致性;顯存中的一部分可以既是portable memory又是mapped memory;在執行CUDA操作前,先調用cudaSetDeviceFlags()(加cudaDeviceMapHost標誌)進行頁鎖定記憶體映射。

 

constant memory:唯讀地址空間;位於顯存,有緩衝加速;64Kb;用於儲存需要頻繁訪問的唯讀參數 ;唯讀;使用_constant_ 關鍵字,定義在所有函數之外;兩種常數儲存空間的使用方法:直接在定義時初始化常數儲存空間;定義一個constant數組,然後使用函數進行賦值;

 

texture memory:唯讀;不是一塊專門的儲存空間,而是牽涉到顯存、兩級紋理緩衝、紋理拾取單元的紋理流水線;資料常以一維、二維或者三維數組的形式儲存在顯存中;緩衝加速;可以聲明大小比常數儲存空間大得多;適合實現映像樹立和尋找表;對大量資料的隨機訪問或非對齊訪問有良好的加速效果;在kernel中訪問紋理儲存空間的操作成為紋理拾取(texture fetching);紋理拾取使用的座標與資料在顯存中的位置可以不同,通過紋理參照系約定二者的映射方式;將顯存中的資料與紋理參照系關聯的操作,稱為將資料與紋理綁定(texture binding);顯存中可以綁定到紋理的資料有:普通線性儲存空間和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對同一Block Storage器操作的順序一致性;顯存中的一部分可以既是portable memory又是mapped memory;在執行CUDA操作前,先調用cudaSetDeviceFlags()(加cudaDeviceMapHost標誌)進行頁鎖定記憶體映射。

 

constant memory:唯讀地址空間;位於顯存,有緩衝加速;64Kb;用於儲存需要頻繁訪問的唯讀參數 ;唯讀;使用_constant_ 關鍵字,定義在所有函數之外;兩種常數儲存空間的使用方法:直接在定義時初始化常數儲存空間;定義一個constant數組,然後使用函數進行賦值;

 

texture memory:唯讀;不是一塊專門的儲存空間,而是牽涉到顯存、兩級紋理緩衝、紋理拾取單元的紋理流水線;資料常以一維、二維或者三維數組的形式儲存在顯存中;緩衝加速;可以聲明大小比常數儲存空間大得多;適合實現映像樹立和尋找表;對大量資料的隨機訪問或非對齊訪問有良好的加速效果;在kernel中訪問紋理儲存空間的操作成為紋理拾取(texture fetching);紋理拾取使用的座標與資料在顯存中的位置可以不同,通過紋理參照系約定二者的映射方式;將顯存中的資料與紋理參照系關聯的操作,稱為將資料與紋理綁定(texture binding);顯存中可以綁定到紋理的資料有:普通線性儲存空間和cuda數組;存在緩衝機制;可以設定濾波模式,定址模式等;

聯繫我們

該頁面正文內容均來源於網絡整理,並不代表阿里雲官方的觀點,該頁面所提到的產品和服務也與阿里云無關,如果該頁面內容對您造成了困擾,歡迎寫郵件給我們,收到郵件我們將在5個工作日內處理。

如果您發現本社區中有涉嫌抄襲的內容,歡迎發送郵件至: info-contact@alibabacloud.com 進行舉報並提供相關證據,工作人員會在 5 個工作天內聯絡您,一經查實,本站將立刻刪除涉嫌侵權內容。

A Free Trial That Lets You Build Big!

Start building with 50+ products and up to 12 months usage for Elastic Compute Service

  • Sales Support

    1 on 1 presale consultation

  • After-Sales Support

    24/7 Technical Support 6 Free Tickets per Quarter Faster Response

  • Alibaba Cloud offers highly flexible support services tailored to meet your exact needs.