日韩性视频-久久久蜜桃-www中文字幕-在线中文字幕av-亚洲欧美一区二区三区四区-撸久久-香蕉视频一区-久久无码精品丰满人妻-国产高潮av-激情福利社-日韩av网址大全-国产精品久久999-日本五十路在线-性欧美在线-久久99精品波多结衣一区-男女午夜免费视频-黑人极品ⅴideos精品欧美棵-人人妻人人澡人人爽精品欧美一区-日韩一区在线看-欧美a级在线免费观看

歡迎訪問 生活随笔!

生活随笔

當前位置: 首頁 > 编程资源 > 编程问答 >内容正文

编程问答

CUDA学习笔记之 CUDA存储器模型

發布時間:2025/3/15 编程问答 22 豆豆
生活随笔 收集整理的這篇文章主要介紹了 CUDA学习笔记之 CUDA存储器模型 小編覺得挺不錯的,現在分享給大家,幫大家做個參考.

CUDA學習筆記之 CUDA存儲器模型

標簽: cuda存儲bindingcache編程api 1223人閱讀 評論(0) 收藏 舉報 分類: CUDA(26) GPU片內:register,shared memory; 板載顯存:local memory,constant memory, texture memory, texture memory,global memory; host 內存: host memory, pinned memory.

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數組;存在緩存機制;可以設定濾波模式,尋址模式等;

與50位技術專家面對面20年技術見證,附贈技術全景圖

總結

以上是生活随笔為你收集整理的CUDA学习笔记之 CUDA存储器模型的全部內容,希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯,歡迎將生活随笔推薦給好友。