GPU存儲器訪問模式_第1頁
GPU存儲器訪問模式_第2頁
GPU存儲器訪問模式_第3頁
GPU存儲器訪問模式_第4頁
GPU存儲器訪問模式_第5頁
已閱讀5頁,還剩19頁未讀, 繼續(xù)免費閱讀

下載本文檔

版權(quán)說明:本文檔由用戶提供并上傳,收益歸屬內(nèi)容提供方,若內(nèi)容存在侵權(quán),請進行舉報或認領(lǐng)

文檔簡介

2012.3.5GPU存儲器訪問模式目錄:

一、GPU簡介

二、GPU存儲器訪問模式

一、GPU簡介GPU(GraphicsProcessingUnit),即圖像處理器,是一種專用圖形渲染設(shè)備。GPU的功能更新很迅速,平均每一年多便有新一代的GPU誕生,運算速度也越來越快。早期GPU主要用于信號成像和圖像處理,GPGPU(GeneralPurposeComputingonGPU)技術(shù)提高了GPU通用計算執(zhí)行效率。GPU實物圖:Teslac2050實物圖

GPU芯片實物圖

GPU的硬件架構(gòu):

右圖是NVIDIA

Geforce8800硬件架構(gòu)圖

GF100架構(gòu)概覽

16個SM(StreamMultiprocessor)每個SM中有32個SP(標量處理器)6個64bitGDDR5存儲器控制器

計算能力2.x的SM(streamingmultiprocessor)的解析圖

?32SP(標量處理器),雙warp-schduler?64KBsharedmemory/L1cache,可以配置為48KBshared/16KBL1或者16KBshared/48KBL1

CUDA

CUDA(ComputingUnifiedDeviceArchitecture)是一種不需要借助于圖形學API進行GPU通用計算的軟硬件架構(gòu)。NVIDIA在2007年推向市場的并行架構(gòu),CUDA并不是一種編程語言,其包括了NVIDIA對于GPGPU的完整解決方案。提供了CUDAC編程語言,類似c語言。CUDA用一種分層的編程模式來組織線程.

SIMT:SIMD的變形

每個block是SIMT,而block間則是MIMDBlock的寬度是可變的,與硬件無關(guān);而通常SIMD編程模型中的向量寬度是固定的一個warp內(nèi)的線程行為更類似SIMDBlock內(nèi)允許存在分支Block內(nèi)的線程間通過sharedmemory和同步進行通信;而SIMD編程模型內(nèi)的向量之間通過寄存器通信,不需要顯式的同步

Gridblockthread

?Kernel以包含大量的線程的Grid形式?存在,每一個線程執(zhí)行相同的程序?Grid由block組成,block內(nèi)的線程間可以高效通信?Block和thread擁有唯一的ID,用于控制thread操作不同的數(shù)據(jù),或者執(zhí)行不同分支Warp?Warp并不存在于抽象的CUDA編程模型,而是由硬件決定的,對性能影響較大?在1.x和2.x架構(gòu)硬件中,一個warp由同一block內(nèi)相鄰的32個線程組成?少數(shù)指令,如vote和ballot,是顯式的以warp為單位執(zhí)行。?同一warp內(nèi)的線程可以認為是“同時”執(zhí)行的,不需要進行同步也能通過sharedmemory進行通信?分支以warp為單位進行?以warp為單位考慮對存儲器的訪問的優(yōu)化,可以獲得更高性能

CUDA開發(fā)環(huán)境:

---支持CUDA的GPU---NVIDIA設(shè)備驅(qū)動器---CUDA開發(fā)工具(可以從NVIDIA下載,用于開發(fā)的工具包)---標準C編譯器CUDA執(zhí)行模型(1)、調(diào)用核程序(有GPU運行的程序)時CPU調(diào)用API將顯卡端程序的代碼傳到GPU(2)、block運行在SM上(3)、thread運行在SP上二、GPU存儲器訪問模式GPU可尋址的存儲器分為:寄存器、全局存儲器、本地存儲器、共享存儲器、常量存儲器以及紋理存儲器。GPU的CUDA硬件整體架構(gòu)圖如下:

數(shù)據(jù)流向:

編址方式---GPU與CPU獨立編址,主機和設(shè)備之間的通信是PCI-E總線通信。---在主機端通過API去分配/復(fù)制/釋放GPU上的存儲空間,類似c的內(nèi)存操作函數(shù)。如:cudaMalloc()函數(shù)用于分配GPU顯存。---GPU不同存儲器統(tǒng)一編址空間,GPU內(nèi)不同存儲器使用相同的編址空間,編譯器可以解析指針指向哪種類型的存儲器---支持64bit尋址空間,最大1TB顯存Register/LocalMemory---線程私有,在kernel程序中不加前綴定義的私有變量---通常情況下被編譯為寄存器,當寄存器資源不足時自動使用LocalMemory/cache作為私有變量---寄存器通常無延遲---計算能力2.x硬件上Local

memory(位于顯存)有cache機制,但是訪問Local

memory同有數(shù)百個周期延遲

SharedMemory---SharedMemory被分為多個組(稱為bank),每個組只有一個32bit端口---相鄰的32bit存儲在相鄰的bank中---Block私有,block線程內(nèi)不同線程通過sharedmemory進行通信---屬于片上存儲器,帶寬較大,延遲很小(無bankconflict時1-2周期)---計算機能力2.x硬件每SM32KBsharedmemory,分為32bank,bank寬度32bit,半速SharedMemorybankconflict

如果同一個half-warp(1.x)或者warp(2.x)訪問了不同端口中的不同數(shù)據(jù)(即同一個bank中的不同字),那么就會發(fā)生bankconflict;---廣播可以避免一部分讀取時的bankconflict,將數(shù)據(jù)廣播給提出讀請求的線程,寫操作只有其中的一個線程寫。---計算機能力2.x硬件有32個bank,訪問64bit型數(shù)據(jù)時不再有bankconflict。

右圖:a、左:以一個字(32位)跨度,線程的一個線性訪問,不存在bank沖突b、中:以一個字(32位)跨度,線程的一個線性訪問,但是其每兩個線程都訪問同一個bank,這就2路bank沖突。c、右:同a,不存在沖突

Globalmemory---通過_device_前綴定義,或者通過API分配---存在于顯存中,訪問延遲大(400-600周期)---受顯存和存儲器控制器設(shè)計影響,按照一定方式(對齊)訪問時才能獲得最大帶寬---計算機能力2.x硬件上有對globalmemory的cache機制對齊1)、圖上部分:地址從128到256的對齊內(nèi)存段,線程順序訪問內(nèi)存,計算能力2.0的設(shè)備中只需一次內(nèi)存操作即可讀完此內(nèi)存段。2)、圖中部分:地址從128到256的對齊內(nèi)存段,線程有部分交叉訪問內(nèi)存,計算能力2.0的設(shè)備中只需一內(nèi)存操作即可讀完此內(nèi)存段。3)、圖下部分:線程訪問起始地址從比128大到超過256的非對齊內(nèi)存段,線程順序訪問內(nèi)存,計算能力2.0的設(shè)備中需兩次內(nèi)存操作即可讀完此內(nèi)存段。

注:warp執(zhí)行指令時,會將warp中多個線程的內(nèi)存訪問集中組織成一次或多次內(nèi)存操作,操作數(shù)依賴于每個線程能訪問的字大小以及內(nèi)存的分配方式。Texturememory---定義紋理參考后與CUDAarray或者Globalmemory綁定---存在于顯存,有texturecache---利用二維三維數(shù)據(jù)的數(shù)據(jù)局部性,訪問相鄰元素性能較好---可以利用GPU中圖形硬件單元進行一些處理,不占用可編程計算性能,計算機能力2.x可寫。---紋理存儲器在每個GPC(圖像處理簇)上有片上紋理緩存---訪問不需要遵循特定的優(yōu)化規(guī)則---尋址的計算由特定的硬件單元實現(xiàn)---紋理緩存提供將輸入的8位或者16位整型映射到[0.0,1.0]或者[-1.0,1.0]的取值空間

紋理存儲器的具體操作:

紋理引用:指定紋理在紋理存儲器中的位置

紋理拾?。涸诮o定紋理引用和紋理坐標后,從紋理上采樣紋理元的過程

紋理坐標:紋理元在紋理上的位置由紋理坐標來描述

紋理尋址:由硬件自動完成

溫馨提示

  • 1. 本站所有資源如無特殊說明,都需要本地電腦安裝OFFICE2007和PDF閱讀器。圖紙軟件為CAD,CAXA,PROE,UG,SolidWorks等.壓縮文件請下載最新的WinRAR軟件解壓。
  • 2. 本站的文檔不包含任何第三方提供的附件圖紙等,如果需要附件,請聯(lián)系上傳者。文件的所有權(quán)益歸上傳用戶所有。
  • 3. 本站RAR壓縮包中若帶圖紙,網(wǎng)頁內(nèi)容里面會有圖紙預(yù)覽,若沒有圖紙預(yù)覽就沒有圖紙。
  • 4. 未經(jīng)權(quán)益所有人同意不得將文件中的內(nèi)容挪作商業(yè)或盈利用途。
  • 5. 人人文庫網(wǎng)僅提供信息存儲空間,僅對用戶上傳內(nèi)容的表現(xiàn)方式做保護處理,對用戶上傳分享的文檔內(nèi)容本身不做任何修改或編輯,并不能對任何下載內(nèi)容負責。
  • 6. 下載文件中如有侵權(quán)或不適當內(nèi)容,請與我們聯(lián)系,我們立即糾正。
  • 7. 本站不保證下載資源的準確性、安全性和完整性, 同時也不承擔用戶因使用這些下載資源對自己和他人造成任何形式的傷害或損失。

評論

0/150

提交評論