存儲(chǔ)器訪問模式_第1頁
存儲(chǔ)器訪問模式_第2頁
存儲(chǔ)器訪問模式_第3頁
存儲(chǔ)器訪問模式_第4頁
存儲(chǔ)器訪問模式_第5頁
已閱讀5頁,還剩18頁未讀, 繼續(xù)免費(fèi)閱讀

VIP免費(fèi)下載

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

文檔簡介

會(huì)計(jì)學(xué)1存儲(chǔ)器訪問模式一、GPU簡介GPU(GraphicsProcessingUnit),即圖像處理器,是一種專用圖形渲染設(shè)備。GPU的功能更新很迅速,平均每一年多便有新一代的GPU誕生,運(yùn)算速度也越來越快。早期GPU主要用于信號(hào)成像和圖像處理,GPGPU(GeneralPurposeComputingonGPU)技術(shù)提高了GPU通用計(jì)算執(zhí)行效率。GPU實(shí)物圖:Teslac2050實(shí)物圖

GPU芯片實(shí)物圖

第1頁/共23頁GPU的硬件架構(gòu):

右圖是NVIDIA

Geforce8800硬件架構(gòu)圖第2頁/共23頁

GF100架構(gòu)概覽

16個(gè)SM(StreamMultiprocessor)每個(gè)SM中有32個(gè)SP(標(biāo)量處理器)6個(gè)64bitGDDR5存儲(chǔ)器控制器第3頁/共23頁

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

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

第4頁/共23頁CUDA

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

SIMT:SIMD的變形

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

Gridblockthread

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

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

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

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

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

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

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

第14頁/共23頁SharedMemory---SharedMemory被分為多個(gè)組(稱為bank),每個(gè)組只有一個(gè)32bit端口---相鄰的32bit存儲(chǔ)在相鄰的bank中---Block私有,block線程內(nèi)不同線程通過sharedmemory進(jìn)行通信---屬于片上存儲(chǔ)器,帶寬較大,延遲很?。obankconflict時(shí)1-2周期)---計(jì)算機(jī)能力2.x硬件每SM32KBsharedmemory,分為32bank,bank寬度32bit,半速第15頁/共23頁SharedMemorybankconflict

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

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

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

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

紋理存儲(chǔ)器的具體操作:

紋理引用:指定紋理在紋理存儲(chǔ)器中的位置

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

紋理坐標(biāo):紋理元在紋理上的位置由紋理坐標(biāo)來描述

紋理尋址:由硬件自動(dòng)完成第20頁/共23頁Constantmem

溫馨提示

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

最新文檔

評(píng)論

0/150

提交評(píng)論