計(jì)算機(jī)組成與系統(tǒng)結(jié)構(gòu)課件：并行體系結(jié)構(gòu)

并行體系結(jié)構(gòu)9.1計(jì)算機(jī)體系結(jié)構(gòu)的并行性9.2計(jì)算機(jī)體系結(jié)構(gòu)的分類9.3互連網(wǎng)絡(luò)9.4向量處理機(jī)9.5多處理器系統(tǒng)9.6圖形處理單元體系結(jié)構(gòu)9.7多計(jì)算機(jī)系統(tǒng)9.8云計(jì)算9.9高性能計(jì)算機(jī)發(fā)展現(xiàn)狀9.10并行處理面臨的挑戰(zhàn)

9.1計(jì)算機(jī)體系結(jié)構(gòu)的并行性

并行計(jì)算是指同時(shí)對(duì)多個(gè)任務(wù)、多條指令或多個(gè)數(shù)據(jù)項(xiàng)進(jìn)行的處理,完成并行計(jì)算的計(jì)算機(jī)系統(tǒng)稱為并行計(jì)算機(jī)系統(tǒng),它是將多個(gè)處理器、多個(gè)計(jì)算機(jī)通過(guò)網(wǎng)絡(luò)以一定的連接方式組織起來(lái)的大規(guī)模系統(tǒng)。

并行計(jì)算機(jī)系統(tǒng)最主要的特性就是并行性。并行性是指計(jì)算機(jī)系統(tǒng)同時(shí)運(yùn)算或同時(shí)操作的特性,它包括同時(shí)性與并發(fā)性兩種含義。同時(shí)性是指兩個(gè)或兩個(gè)以上事件在同一時(shí)刻發(fā)生,并發(fā)性是指兩個(gè)或兩個(gè)以上事件在同一時(shí)間間隔內(nèi)發(fā)生。并行機(jī)制可以加速計(jì)算機(jī)系統(tǒng)的運(yùn)行速度。

在計(jì)算機(jī)系統(tǒng)中,實(shí)現(xiàn)并行機(jī)制的途徑有多種,可以歸納為以下四類:

(1)時(shí)間重疊:運(yùn)用的是時(shí)間并行技術(shù),它通過(guò)讓多個(gè)處理任務(wù)或子任務(wù)同時(shí)使用系統(tǒng)中的不同功能部件,使系統(tǒng)處理任務(wù)的吞吐量增大,從而達(dá)到使系統(tǒng)運(yùn)行速度提高的目的。

(2)資源重復(fù):運(yùn)用的是空間并行技術(shù),它通過(guò)大量重復(fù)設(shè)置硬件資源,使多個(gè)處理任務(wù)或子任務(wù)同時(shí)使用系統(tǒng)中的多個(gè)相同功能的部件,使系統(tǒng)處理任務(wù)的吞吐量增大,從而達(dá)到使系統(tǒng)運(yùn)行速度提高的目的。

(3)時(shí)間重疊+資源重復(fù):同時(shí)運(yùn)用時(shí)間并行和空間并行技術(shù),這已成為當(dāng)前并行機(jī)制的主流。

(4)資源共享:與前三種硬件方式不同,這是一種軟件方式,通過(guò)操作系統(tǒng)的調(diào)度使多個(gè)任務(wù)按一定規(guī)則(如時(shí)間片)輪流使用同一設(shè)備。資源共享既可以降低成本,提高設(shè)備利用率,又可以實(shí)現(xiàn)多任務(wù)分時(shí)并行處理。

在計(jì)算機(jī)系統(tǒng)中,可以在不同層次引入并行機(jī)制,如圖9.1所示。當(dāng)多個(gè)CPU或者處理元件緊密連在一起時(shí),它們之間具有高帶寬和低時(shí)延,而且是親密計(jì)算,通常稱之為緊耦合(TightlyCoupled);相反,當(dāng)它們距離較遠(yuǎn)時(shí),具有低帶寬、高時(shí)延,而且是遠(yuǎn)程計(jì)算,通常稱之為松耦合(LooselyCoupled)。從片內(nèi)并行、協(xié)處理器、多處理器、多計(jì)算機(jī)到云計(jì)算,隨著計(jì)算機(jī)體系結(jié)構(gòu)并行層次的上移,系統(tǒng)從最底層緊耦合的系統(tǒng)開(kāi)始逐步轉(zhuǎn)變到高層松耦合的系統(tǒng)。圖9.1計(jì)算機(jī)體系結(jié)構(gòu)的并行層次

1.片內(nèi)并行

片內(nèi)并行是在計(jì)算機(jī)體系結(jié)構(gòu)的最底層———芯片內(nèi)部引入的一種并行機(jī)制,即并行行為都發(fā)生在一個(gè)單獨(dú)的芯片內(nèi)部。它實(shí)現(xiàn)加速的方法是使芯片在同一時(shí)間內(nèi)完成更多工作,從而增加芯片吞吐量。片內(nèi)并行的第一種形式是指令級(jí)并行,允許多條指令在片內(nèi)流水線的不同功能單元上并行執(zhí)行。第二種形式是芯片多線程,在這種并行中,CPU可以在多個(gè)線程之間來(lái)回切換,產(chǎn)生虛擬多處理器。第三種形式是單片多處理器(多核CPU),即在同一個(gè)芯片中設(shè)置了兩個(gè)或者更多個(gè)處理器內(nèi)核,并且允許它們同時(shí)運(yùn)行。

2.協(xié)處理器

協(xié)處理器(Coprocessor)是為減輕主處理器負(fù)擔(dān)、協(xié)助主處理器完成特定工作的專用處理器。主處理器和協(xié)處理器并行工作可使計(jì)算機(jī)速度得到提高。

指令級(jí)并行對(duì)提高速度有幫助,但流水線和超標(biāo)量體系結(jié)構(gòu)對(duì)速度的提高很難超過(guò)10倍以上。如果要將計(jì)算機(jī)系統(tǒng)性能提高百倍、千倍甚至百萬(wàn)倍,唯一的辦法就是使用多個(gè)甚至成千上萬(wàn)個(gè)CPU,讓它們連接成一個(gè)大的系統(tǒng),一起并行高效地工作。由多個(gè)CPU構(gòu)成的系統(tǒng)分為多處理器系統(tǒng)和多計(jì)算機(jī)系統(tǒng)。這兩種設(shè)計(jì)的核心是CPU間的相互通信,主要區(qū)別是它們是否有共享的內(nèi)存,這種區(qū)別影響著并行計(jì)算機(jī)系統(tǒng)的設(shè)計(jì)、構(gòu)建、編程、規(guī)模和價(jià)格。

3.多處理器

所有的CPU共享公共內(nèi)存的并行計(jì)算機(jī)稱為多處理器系統(tǒng),如圖9.2(a)所示。運(yùn)行在多處理器上的所有進(jìn)程能夠共享映射到公共內(nèi)存的單一虛擬地址空間。任何進(jìn)程都能通過(guò)執(zhí)行Load或者Store指令來(lái)讀或?qū)懸粋€(gè)內(nèi)存字,其余工作由硬件來(lái)完成。采用一個(gè)進(jìn)程先把數(shù)據(jù)寫(xiě)入內(nèi)存然后由另一個(gè)進(jìn)程讀出的方式,兩個(gè)進(jìn)程之間可以進(jìn)行通信。圖9.2并行體系結(jié)構(gòu)

4.多計(jì)算機(jī)

多計(jì)算機(jī)系統(tǒng)是由具有大量CPU但不共享公共內(nèi)存的系統(tǒng)構(gòu)成的。在多計(jì)算機(jī)系統(tǒng)中,每個(gè)CPU都有私有本地內(nèi)存,私有內(nèi)存只能供CPU自己通過(guò)執(zhí)行Load和Store指令來(lái)使用,其他CPU則不能直接訪問(wèn),這種體系結(jié)構(gòu)也被稱為分布式內(nèi)存系統(tǒng)(DistributedMemorySystem,DMS),如圖9.2(b)所示。當(dāng)處理器Pi發(fā)現(xiàn)Pj有它需要的數(shù)據(jù)時(shí),它給Pj發(fā)送一條請(qǐng)求數(shù)據(jù)的消息,然后Pi進(jìn)入阻塞(即等待),直到請(qǐng)求被響應(yīng)。當(dāng)消息到達(dá)Pj后,Pj的軟件將分析該消息并把響應(yīng)消息和需要的數(shù)據(jù)發(fā)送給Pi。當(dāng)響應(yīng)消息到達(dá)Pi后,Pi軟件將解除阻塞、接收數(shù)據(jù)并繼續(xù)執(zhí)行。所以,進(jìn)程間通信通常使用Send和Receive這樣的軟件原語(yǔ)實(shí)現(xiàn)。

5.云計(jì)算

到目前為止,使具有不同計(jì)算機(jī)操作系統(tǒng)、數(shù)據(jù)庫(kù)和協(xié)議的不同組織一起工作,進(jìn)而共享資源和數(shù)據(jù),還是非常困難的。然而,不斷增長(zhǎng)的對(duì)大規(guī)模組織間協(xié)作的需求引領(lǐng)了新的系統(tǒng)和技術(shù)的開(kāi)發(fā)。云計(jì)算(CloudComputing)是網(wǎng)格計(jì)算(GridComputing)、分布計(jì)算(DistributedComputing)、并行計(jì)算(ParallelComputing)、效用計(jì)算(UtilityComputing)、網(wǎng)絡(luò)存儲(chǔ)技術(shù)(NetworkStorageTechnologies)、虛擬化(Virtualization)、負(fù)載均衡(LoadBalance)等傳統(tǒng)計(jì)算機(jī)和網(wǎng)絡(luò)技術(shù)發(fā)展融合的產(chǎn)物。云計(jì)算的核心思想是將大量用網(wǎng)絡(luò)連接的計(jì)算資源統(tǒng)一管理和調(diào)度,構(gòu)成一個(gè)計(jì)算資源池,對(duì)用戶按需服務(wù)。提供資源的網(wǎng)絡(luò)被稱為云。

云計(jì)算是一種基于互聯(lián)網(wǎng)的計(jì)算方式,通過(guò)這種方式,云中的共享軟硬件資源和信息在使用者看來(lái)是可以無(wú)限擴(kuò)展的,并且可以隨時(shí)獲取,按需使用,按使用付費(fèi)。許多云計(jì)算部署依賴于計(jì)算機(jī)集群,整個(gè)運(yùn)行方式很像電網(wǎng)。云計(jì)算系統(tǒng)可以被視為非常大的、國(guó)際間的、松耦合的、異構(gòu)的集群集合。

9.2計(jì)算機(jī)體系結(jié)構(gòu)的分類

Flynn于1966年提出了一種今天仍有價(jià)值的對(duì)所有計(jì)算機(jī)進(jìn)行分類的簡(jiǎn)單模型,這種分類模型可以為計(jì)算機(jī)系統(tǒng)設(shè)計(jì)制訂一個(gè)框架,這就是1.4.1節(jié)中介紹的Flynn分類法。根據(jù)被調(diào)用的數(shù)據(jù)流和指令流的并行度,Flynn分類法將計(jì)算機(jī)歸為以下四類:

(1)單指令流單數(shù)據(jù)流(SingleInstruction-streamSingleData-stream,SISD)。這是一種單處理器系統(tǒng),傳統(tǒng)的馮·諾依曼計(jì)算機(jī)就是SISD系統(tǒng)。它只有一個(gè)指令流和一個(gè)數(shù)據(jù)流,一個(gè)時(shí)刻只能做一件事情。

(2)單指令流多數(shù)據(jù)流(SingleInstruction-streamMultipleData-stream,SIMD)。SIMD計(jì)算機(jī)中,同一條指令被多個(gè)使用不同數(shù)據(jù)流的多處理器執(zhí)行。

(3)多指令流單數(shù)據(jù)流(MultipleInstruction-streamSingleData-stream,MISD)。MISD計(jì)算機(jī)中,多條指令同時(shí)在同一數(shù)據(jù)上進(jìn)行操作。至今還沒(méi)有這類的商用機(jī)器。

(4)多指令流多數(shù)據(jù)流(MultipleInstruction-streamMultipleData-stream,MIMD)。這是一種同時(shí)有多個(gè)CPU執(zhí)行不同操作的計(jì)算機(jī)系統(tǒng),系統(tǒng)中每個(gè)處理器獲取自己的指令并對(duì)自己的數(shù)據(jù)進(jìn)行操作處理。

Flynn是一種粗略的分類,隨著并行體系結(jié)構(gòu)的發(fā)展,有些計(jì)算機(jī)成為上述多種類型的混合體,有些計(jì)算機(jī)成為上述某種類型的子類?，F(xiàn)在普遍采用的更為細(xì)致的分類,如圖9.3所示。圖9.3計(jì)算機(jī)分類

多計(jì)算機(jī)系統(tǒng)又可以分成以下兩大類:

(1)大規(guī)模并行處理(MassivelyParallelProcessing,MPP)系統(tǒng)。

(2)集群(Cluster)。它是由普通的PC或者工作站組成的,它們可能被放置在一個(gè)大的機(jī)架或地域上,相互之間通過(guò)現(xiàn)成的商用網(wǎng)絡(luò)連接。

MIMD靈活性強(qiáng),在必要的軟件和硬件支持下,MIMD既能作為單用戶多處理器為單一應(yīng)用程序提供高性能,又可作為多道程序多處理器系統(tǒng)同時(shí)運(yùn)行多個(gè)任務(wù),還可以提供結(jié)合這兩種任務(wù)的應(yīng)用。MIMD能夠充分利用現(xiàn)有微處理器的性價(jià)比優(yōu)勢(shì),當(dāng)今幾乎所有商用多處理器系統(tǒng)使用的微處理器與在工作站、單處理器服務(wù)器中所使用的微處理器都是相同的。從圖9.3的分類中可以清楚地看到,MIMD計(jì)算機(jī)是高性能計(jì)算機(jī)發(fā)展的主流。

9.3互連網(wǎng)絡(luò)

9.3.1基本概念互連網(wǎng)絡(luò)(InterconnectionNetwork)是一種由開(kāi)關(guān)元件按照一定的拓?fù)浣Y(jié)構(gòu)和控制方式構(gòu)成的網(wǎng)絡(luò),用于實(shí)現(xiàn)計(jì)算機(jī)系統(tǒng)中部件之間、處理器之間、部件與處理器之間,甚至計(jì)算機(jī)之間的相互連接。它的主要構(gòu)成元素包括終端節(jié)點(diǎn)、鏈路、連接節(jié)點(diǎn)。

設(shè)備(Device)是計(jì)算機(jī)內(nèi)部的一個(gè)部件或一組部件(如CPU、存儲(chǔ)器),也可以是計(jì)算機(jī)或計(jì)算機(jī)系統(tǒng)。終端節(jié)點(diǎn)(EndNode)由設(shè)備及相關(guān)的軟硬件接口組成;鏈路(Link)是節(jié)點(diǎn)與互連網(wǎng)絡(luò)及互連網(wǎng)絡(luò)內(nèi)部的連接線,它可以是有向或無(wú)向的;連接節(jié)點(diǎn)(ConnectedNode)用于接入節(jié)點(diǎn)設(shè)備和連接不同的鏈路,主要由高速開(kāi)關(guān)構(gòu)成?；ミB網(wǎng)絡(luò)也稱為通信子網(wǎng)(CommunicationSubnet)或通信子系統(tǒng)(CommunicationSubsystem),多個(gè)網(wǎng)絡(luò)的相互連接稱為網(wǎng)絡(luò)互連(Internetworking)。

根據(jù)連接的設(shè)備數(shù)和設(shè)備的接近程度,可以將互連網(wǎng)絡(luò)分為以下四類:

(1)片上網(wǎng)(On-ChipNetwork,OCN)。OCN也稱為NoC(Network-on-Chip),用于連接微體系結(jié)構(gòu)中的功能單元以及芯片或多芯片模塊內(nèi)的處理器和IP核(IntellectualPropertycore,知識(shí)產(chǎn)權(quán)核)。

(2)系統(tǒng)/存儲(chǔ)區(qū)域網(wǎng)(System/StorageAreaNetwork,SAN)。系統(tǒng)區(qū)域網(wǎng)用于多處理器和多計(jì)算機(jī)系統(tǒng)內(nèi)的處理器之間、計(jì)算機(jī)之間的互連;存儲(chǔ)區(qū)域網(wǎng)用于處理器與存儲(chǔ)器之間、計(jì)算機(jī)系統(tǒng)與存儲(chǔ)系統(tǒng)之間的互連,也用于在服務(wù)器和數(shù)據(jù)中心環(huán)境中的存儲(chǔ)設(shè)備與I/O設(shè)備之間的連接。

(3)局域網(wǎng)(LocalAreaNetwork,LAN)。LAN用于互連分布在機(jī)房、大廈或校園環(huán)境中的自治計(jì)算機(jī)系統(tǒng)。

(4)廣域網(wǎng)(WideAreaNetwork,WAN)。WAN連接分布在全球的計(jì)算機(jī)系統(tǒng),需要互連網(wǎng)絡(luò)的支持。WAN可以連接相距幾千公里的上百萬(wàn)個(gè)計(jì)算機(jī)。ATM是WAN的一個(gè)示例。

計(jì)算機(jī)體系結(jié)構(gòu)關(guān)注互連網(wǎng)絡(luò)有以下幾個(gè)原因:

(1)互連網(wǎng)絡(luò)除了提供連通性外,還經(jīng)常用于在多個(gè)層次上連接計(jì)算機(jī)部件,包括在處理器微體系結(jié)構(gòu)(Microarchitecture)中。

(2)網(wǎng)絡(luò)一直用于大型機(jī),而今天這種設(shè)計(jì)在個(gè)人計(jì)算機(jī)中也得到了很好的應(yīng)用,它提供的高通信帶寬使計(jì)算能力和存儲(chǔ)容量得以提升。

(3)拓?fù)浣Y(jié)構(gòu)與其他網(wǎng)絡(luò)設(shè)計(jì)參數(shù)之間存在著敏感的關(guān)系,特別是當(dāng)終端節(jié)點(diǎn)數(shù)量非常大(如BlueGene/L巨型計(jì)算機(jī)中的64K個(gè)節(jié)點(diǎn))或當(dāng)?shù)却龝r(shí)間(Latency)成為關(guān)鍵指標(biāo)(如在多核處理器芯片中)時(shí)。

(4)拓?fù)浣Y(jié)構(gòu)也極大地影響著網(wǎng)絡(luò)的實(shí)現(xiàn)成本和整個(gè)計(jì)算機(jī)系統(tǒng)的性能價(jià)格比。

所以,為了更有效地設(shè)計(jì)和評(píng)估計(jì)算機(jī)系統(tǒng),計(jì)算機(jī)設(shè)計(jì)師應(yīng)該了解互連問(wèn)題及解決方法?；ミB網(wǎng)絡(luò)設(shè)計(jì)的宗旨是不應(yīng)成為系統(tǒng)性能和成本的瓶頸。

9.3.2互連網(wǎng)絡(luò)的表示及性能參數(shù)

1.圖

互連網(wǎng)絡(luò)可用圖表示,圖是由有向邊或無(wú)向邊連接有限個(gè)節(jié)點(diǎn)組成的,通常認(rèn)為圖中所有的節(jié)點(diǎn)或所有的邊可以是同質(zhì)的(Homogeneity)。邊表示節(jié)點(diǎn)間互連的鏈路,在節(jié)點(diǎn)上的節(jié)點(diǎn)設(shè)備通過(guò)邊進(jìn)行信息交換。每個(gè)節(jié)點(diǎn)都有邊與之相連,數(shù)學(xué)上將與節(jié)點(diǎn)相連的邊數(shù)稱為節(jié)點(diǎn)度(Degree),進(jìn)入節(jié)點(diǎn)的邊數(shù)叫入度(扇入,Fanin),從節(jié)點(diǎn)出來(lái)的邊數(shù)則叫出度(扇出,Fanout)。一般來(lái)說(shuō),扇入/扇出越大,路由選擇能力越強(qiáng),容錯(cuò)能力也越強(qiáng)。容錯(cuò)是指當(dāng)某條鏈路失效時(shí)可以繞過(guò)這條鏈路繼續(xù)保持系統(tǒng)正常工作。

互連網(wǎng)絡(luò)的一個(gè)重要特性是傳輸延遲,可以用圖的直徑(Diameter)來(lái)表示。如果用兩個(gè)節(jié)點(diǎn)之間的邊數(shù)來(lái)表示兩個(gè)節(jié)點(diǎn)之間的距離,那么圖的直徑就是圖中相距最遠(yuǎn)的兩個(gè)節(jié)點(diǎn)之間的距離。圖的直徑直接關(guān)系到CPU與CPU之間、CPU與內(nèi)存之間,甚至計(jì)算機(jī)與計(jì)算機(jī)之間交換信息包時(shí)的最大延遲。因?yàn)樾畔⑼ㄟ^(guò)每條鏈路都要花費(fèi)一定的時(shí)間,所以圖的直徑越小,最壞情況下的互連網(wǎng)絡(luò)性能就越好。兩個(gè)節(jié)點(diǎn)之間的平均距離也很重要,它關(guān)系到信息包的平均傳遞時(shí)間。

一個(gè)信息經(jīng)過(guò)互連網(wǎng)絡(luò)到達(dá)接收方造成的總時(shí)間延遲為

其中,頻帶寬度(BandWidth,簡(jiǎn)稱頻寬)是互連網(wǎng)絡(luò)傳輸信息的最大速率;傳輸時(shí)間(TransmissionTime)等于信息包長(zhǎng)度除以頻帶寬度;飛行時(shí)間(TimeofFlight)是第一位信息到達(dá)接收方所花費(fèi)的時(shí)間;傳輸延遲(TransportLatency)等于飛行時(shí)間與傳輸時(shí)間之和;發(fā)送方開(kāi)銷(SenderOverhead)是處理器把消息放到互連網(wǎng)絡(luò)的時(shí)間;接收方開(kāi)銷(ReceiverOverhead)是處理器把消息從網(wǎng)絡(luò)取出來(lái)的時(shí)間。

按照?qǐng)D的維數(shù)(Dimension)可以對(duì)互連網(wǎng)絡(luò)進(jìn)行分類。圖的維數(shù)定義為源節(jié)點(diǎn)和目的節(jié)點(diǎn)之間可供選擇的路徑數(shù)量。如果只有一種選擇,圖和網(wǎng)絡(luò)就是一維的;如果有兩種選擇,圖和網(wǎng)絡(luò)就是二維的;如果有三種選擇,圖和網(wǎng)絡(luò)就是三維的,以此類推,可以得到n維網(wǎng)絡(luò)。

2.互連函數(shù)

互連網(wǎng)絡(luò)也可以用互連函數(shù)表示,互連函數(shù)定義為互連網(wǎng)絡(luò)輸入/輸出端口地址的一對(duì)一映射(Bijection),即

其中,x為自變量,是互連網(wǎng)絡(luò)輸入端口編號(hào)(即地址)的n位二進(jìn)制編碼;y為因變量,是互連網(wǎng)絡(luò)輸出端口編號(hào)的n位二進(jìn)制編碼;f()表示互連網(wǎng)絡(luò)輸入與輸出的對(duì)應(yīng)關(guān)系,它可以實(shí)現(xiàn)對(duì)x編碼進(jìn)行排列、組合、移位、取反等操作。一個(gè)互連網(wǎng)絡(luò)的連接特征可以對(duì)應(yīng)多個(gè)互連函數(shù)。

互連函數(shù)有以下幾種表示方法:

(1)函數(shù)表示法:用x表示輸入端變量,用f(x)表示互連函數(shù),如公式(9.2)。

(2)表格表示法:適用于規(guī)則和不規(guī)則連接。例如:

(3)循環(huán)表示法:將有輸入/輸出連接且編號(hào)能夠銜接并循環(huán)的一組連線的節(jié)點(diǎn)放在一個(gè)括號(hào)內(nèi)來(lái)表示輸入端口與輸出端口的映射關(guān)系,如存在輸入端口→輸出端口的連接為0→2、2→4、4→6、6→0,則有(0246)。

(4)圖形表示法:用連線表示輸入/輸出端口的映射關(guān)系。

典型的網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)有總線結(jié)構(gòu)、集中式交換網(wǎng)絡(luò)結(jié)構(gòu)和分布式交換網(wǎng)絡(luò)結(jié)構(gòu)。對(duì)共享總線的一種替代是交換網(wǎng)絡(luò),這是一種利用若干小交換節(jié)點(diǎn)并以不同方式連接而形成的多條點(diǎn)對(duì)點(diǎn)鏈路構(gòu)成的互連網(wǎng)絡(luò)。根據(jù)連接對(duì)象距離的遠(yuǎn)近,交換網(wǎng)絡(luò)分為集中式和分布式兩類。

9.3.3集中式交換網(wǎng)絡(luò)

典型的集中式交換網(wǎng)絡(luò)有交叉開(kāi)關(guān)網(wǎng)絡(luò)和多級(jí)互連網(wǎng)絡(luò)。

1.交叉開(kāi)關(guān)網(wǎng)絡(luò)

交叉開(kāi)關(guān)(CrossbarSwitch)技術(shù)來(lái)源于電話網(wǎng)絡(luò)中采用的電路交換技術(shù)。電路交換(CircuitSwitching)專用于在源節(jié)點(diǎn)(Node)和目的節(jié)點(diǎn)之間建立通信通道(CommunicationPath),通道是在節(jié)點(diǎn)間的一組順序連接的物理鏈路(PhysicalLink)。

利用電路交換技術(shù)控制若干開(kāi)關(guān)組成的開(kāi)關(guān)連接陣列,就構(gòu)成了交叉開(kāi)關(guān)網(wǎng)絡(luò),它可以依照控制策略將網(wǎng)絡(luò)的任意輸入與輸出連接起來(lái)。圖9.4是一個(gè)4×4交叉開(kāi)關(guān)網(wǎng)絡(luò),交叉陣列中行線與列線的交叉點(diǎn)是一個(gè)交叉開(kāi)關(guān),即一個(gè)小的交換節(jié)點(diǎn),它的電路狀態(tài)是打開(kāi)或閉合,用黑圈表示閉合的交叉開(kāi)關(guān),用白圈表示打開(kāi)的交叉開(kāi)關(guān)。圖9.44×4交叉開(kāi)關(guān)

例9.1在UMA對(duì)稱多處理器系統(tǒng)中,假設(shè)有4個(gè)CPU作為互連網(wǎng)絡(luò)入端節(jié)點(diǎn),4個(gè)內(nèi)存模塊作為出端節(jié)點(diǎn),當(dāng)將它們以圖9.4所示的交叉開(kāi)關(guān)網(wǎng)絡(luò)進(jìn)行連接時(shí),試分析CPU與內(nèi)存模塊的連接關(guān)系。

解因?yàn)镃PU2與內(nèi)存模塊0之間的連線交叉點(diǎn)為閉合的交叉開(kāi)關(guān)(黑圈),可表示為

所以CPU2與內(nèi)存模塊0之間建立了連接。

交叉開(kāi)關(guān)網(wǎng)絡(luò)的優(yōu)勢(shì):

①它是無(wú)阻塞網(wǎng)絡(luò),即節(jié)點(diǎn)之間不會(huì)因?yàn)槟承┙徊纥c(diǎn)或者鏈路被占用而無(wú)法建立連接(僅在有兩個(gè)以上節(jié)點(diǎn)試圖向同一個(gè)目的節(jié)點(diǎn)發(fā)信息時(shí)才可能在目的節(jié)點(diǎn)接收鏈路上發(fā)生阻塞);

②建立連接時(shí)不需要事先規(guī)劃。事實(shí)上,開(kāi)關(guān)網(wǎng)絡(luò)正在替代總線成為計(jì)算機(jī)之間、I/O設(shè)備之間、電路板之間、芯片之間甚至芯片內(nèi)部模塊之間的常規(guī)通信手段。

交叉開(kāi)關(guān)網(wǎng)絡(luò)的劣勢(shì):它的復(fù)雜度隨網(wǎng)絡(luò)端口數(shù)量以平方級(jí)增長(zhǎng),即交叉開(kāi)關(guān)的數(shù)量達(dá)到n×k個(gè)。對(duì)于中等規(guī)模的系統(tǒng),交叉開(kāi)關(guān)設(shè)計(jì)是可行的,如SunFireE25K就是采用這樣的設(shè)計(jì)。但如果n=k=1000,那么就需要100萬(wàn)個(gè)交叉開(kāi)關(guān),這樣大規(guī)模的交叉開(kāi)關(guān)網(wǎng)絡(luò)是很難實(shí)現(xiàn)的。

對(duì)交叉開(kāi)關(guān)網(wǎng)絡(luò)的一種改進(jìn)是通過(guò)用多個(gè)較小規(guī)模的交叉開(kāi)關(guān)模塊串聯(lián)和并聯(lián)構(gòu)成多級(jí)交叉開(kāi)關(guān)網(wǎng)絡(luò),以取代單級(jí)的大規(guī)模交叉開(kāi)關(guān)。圖9.5是用4×4的交叉開(kāi)關(guān)模塊組成16×16的兩級(jí)交叉開(kāi)關(guān)網(wǎng)絡(luò),其設(shè)備量減少為單級(jí)16×16的一半。這實(shí)際上是用4×4的交叉開(kāi)關(guān)模塊構(gòu)成42×42的交叉開(kāi)關(guān)網(wǎng)絡(luò),其中,指數(shù)2為互連網(wǎng)絡(luò)的級(jí)數(shù)。若互連網(wǎng)絡(luò)的入端數(shù)和出端數(shù)不同,可使用a×b的交叉開(kāi)關(guān)模塊,使a中任一輸入端與b中任一輸出端相連。用n級(jí)a×b交叉開(kāi)關(guān)模塊,可以組成一個(gè)an×bn的開(kāi)關(guān)網(wǎng)絡(luò),稱作Delta網(wǎng),它在Patel(1981年)多處理機(jī)中采用。圖9.5用4×4交叉開(kāi)關(guān)模塊構(gòu)成16×16的兩級(jí)交叉開(kāi)關(guān)網(wǎng)絡(luò)

2.多級(jí)互連網(wǎng)絡(luò)

另一種組織與控制更為有效的交換網(wǎng)絡(luò)是基于a×b交換開(kāi)關(guān)構(gòu)造而成的。2×2交換開(kāi)關(guān)是一種最常用的二元開(kāi)關(guān),如圖9.6(a)所示,它有兩個(gè)輸入和兩個(gè)輸出,從任意輸入線到達(dá)的消息都可以交換到任意的輸出線上。

因?yàn)槊總€(gè)輸入可與一個(gè)或兩個(gè)輸出相連,所以在輸出端必須避免發(fā)生沖突。一對(duì)一和一對(duì)多映射是容許的,但不容許有多對(duì)一映射。只容許一對(duì)一映射時(shí),稱為置換連接,并稱此開(kāi)關(guān)為2×2交叉開(kāi)關(guān)。具有直通和交換兩種功能的交換開(kāi)關(guān)稱為兩功能開(kāi)關(guān),用一位控制信號(hào)控制。具有所有四種功能的交換開(kāi)關(guān)稱為四功能開(kāi)關(guān),用兩位控制信號(hào)控制。交換開(kāi)關(guān)的四種狀態(tài)(直通Through,交叉Cross,上播UpperBroadcast,下播LowerBroadcast)如圖9.6(c)所示,通過(guò)加載控制信號(hào)進(jìn)行選擇。圖9.62×2的交換開(kāi)關(guān)

多級(jí)互連網(wǎng)絡(luò)設(shè)計(jì)的關(guān)鍵是:

(1)選擇何種交換開(kāi)關(guān)。

(2)交換開(kāi)關(guān)之間采用何種拓?fù)溥B接。

(3)對(duì)交換開(kāi)關(guān)采用何種控制方式。

最常選用的交換開(kāi)關(guān)是2×2交換開(kāi)關(guān),因?yàn)樗Y(jié)構(gòu)簡(jiǎn)單,容易控制,成本低。級(jí)間連接常用的拓?fù)溆腥煜?、蝶形、縱橫交叉、立方體連接等。對(duì)交換開(kāi)關(guān)的控制一般有以下四種方式:

(1)整體控制:所有交換開(kāi)關(guān)使用同一個(gè)控制信號(hào)控制。

(2)級(jí)控制:同一級(jí)交換開(kāi)關(guān)使用同一個(gè)控制信號(hào)控制。

(3)單元級(jí)控制:每個(gè)交換開(kāi)關(guān)分別控制。

(4)部分級(jí)控制:同一級(jí)中部分交換開(kāi)關(guān)使用同一個(gè)控制信號(hào)控制。

Omega網(wǎng)絡(luò)(或稱Ω網(wǎng)絡(luò))是一種只提供必要服務(wù)且經(jīng)濟(jì)的多級(jí)互連網(wǎng)絡(luò),如圖9.7所示。圖9.7Omega網(wǎng)絡(luò)

例9.2假設(shè)CPU011需要從內(nèi)存模塊110中讀取一個(gè)字,試分析Omega網(wǎng)絡(luò)為此提供的連接路徑。

解我們用圖9.8來(lái)說(shuō)明Omega網(wǎng)絡(luò)的尋徑過(guò)程。圖9.8例9.2Omega網(wǎng)絡(luò)尋徑示意圖

第二種解決方案是在原有的交換網(wǎng)中再插入期望的交換網(wǎng),在這種情況下,通過(guò)中間級(jí)開(kāi)關(guān)提供的足夠多的選擇路徑,來(lái)允許在第一級(jí)和最后一級(jí)交換開(kāi)關(guān)之間建立無(wú)沖突路徑。最著名的例子就是Clos網(wǎng)絡(luò),如圖9.9所示的是3級(jí)Clos網(wǎng),它是嚴(yán)格無(wú)阻塞網(wǎng)。為了將所有的交換開(kāi)關(guān)節(jié)點(diǎn)規(guī)模降低到2×2,三級(jí)Clos拓?fù)浣Y(jié)構(gòu)的多通道特性被遞歸地應(yīng)用于中間交換級(jí),最終獲得了2×lbN-1級(jí)的Bene?拓?fù)浣Y(jié)構(gòu),這是一種路徑可重規(guī)劃且無(wú)阻塞的網(wǎng)絡(luò)結(jié)構(gòu)。Myrinet網(wǎng)(1994年至今)采用的就是Clos網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)。圖9.93級(jí)Clos網(wǎng)絡(luò)

圖9.10說(shuō)明了兩種Bene?拓?fù)浣Y(jié)構(gòu)。圖9.10(a)所示的是16端口的Bene?拓?fù)浣Y(jié)構(gòu),它的中間級(jí)開(kāi)關(guān)(在虛線框內(nèi))是由一個(gè)Clos網(wǎng)絡(luò)實(shí)現(xiàn)的,而該Clos網(wǎng)絡(luò)的中間級(jí)開(kāi)關(guān)(在點(diǎn)畫(huà)線框內(nèi))又是一個(gè)Clos網(wǎng)絡(luò),如此下去,直到各交換開(kāi)關(guān)只使用2×2交換開(kāi)關(guān)時(shí)Bene?網(wǎng)絡(luò)就生成了,即Bene?網(wǎng)絡(luò)中除第一級(jí)和最后一級(jí)之外的所有中間交換級(jí)由Clos網(wǎng)絡(luò)遞歸生成。

到目前為止所描述的MIN只有單向的網(wǎng)絡(luò)鏈路,而雙向網(wǎng)絡(luò)鏈路很容易通過(guò)簡(jiǎn)單地對(duì)折諸如Clos和Bene?這樣的對(duì)稱網(wǎng)絡(luò)來(lái)獲得。將兩個(gè)方向的單向鏈路重疊使用可獲得雙向鏈路,再讓2組交換開(kāi)關(guān)并聯(lián)工作,可獲得4×4的交換開(kāi)關(guān)。圖9.10(b)所示為最終形成的對(duì)折的Bene?拓?fù)浣Y(jié)構(gòu),這是一個(gè)使用4×4交換開(kāi)關(guān)的雙向Bene?網(wǎng)絡(luò),終端節(jié)點(diǎn)連到了Bene?網(wǎng)絡(luò)(單向的)最內(nèi)部的交換開(kāi)關(guān)。在網(wǎng)絡(luò)另一邊的端口是開(kāi)放的,這可以使網(wǎng)絡(luò)進(jìn)一步擴(kuò)展到更大規(guī)模。

這種類型的網(wǎng)絡(luò)被稱為雙向多級(jí)互連網(wǎng)絡(luò)(BidirectionalMultistageInterconnectionNetwork)。在這類網(wǎng)絡(luò)所具有的許多有用性質(zhì)中,最重要的是它們的模塊化和開(kāi)拓通信節(jié)點(diǎn)的能力,這些通信節(jié)點(diǎn)保存著跳過(guò)網(wǎng)絡(luò)各交換級(jí)的信息(消息)包。這類網(wǎng)絡(luò)的規(guī)則性減少了路由的復(fù)雜性,多通道性使通信量更均衡地分布在網(wǎng)絡(luò)資源上,且使網(wǎng)絡(luò)具有了容錯(cuò)能力。

圖9.10兩種Bene?拓?fù)浣Y(jié)構(gòu)

例9.3在Omega網(wǎng)絡(luò)中,級(jí)間連接采用了完全相同的全混洗連接,如圖9.11所示,試寫(xiě)出其互連函數(shù)。

解設(shè)輸入端節(jié)點(diǎn)編號(hào)與輸出端節(jié)點(diǎn)編號(hào)用二進(jìn)制編碼xn-1xn-2…x1x0表示,則輸入與輸出節(jié)點(diǎn)全混洗連接關(guān)系可用如下互連函數(shù)表示:

這表示被連接的一對(duì)輸入/輸出節(jié)點(diǎn)的編碼具有如此關(guān)系:將輸入端節(jié)點(diǎn)編號(hào)的二進(jìn)制編碼循環(huán)左移一位就得到與之相連的輸出端節(jié)點(diǎn)編號(hào)的二進(jìn)制編碼。圖9.11全混洗連接

(2)確定圖9.10(b)的級(jí)間互連函數(shù)。

圖9.10(b)的級(jí)間連接采用的是蝶形連接,其中第2級(jí)的級(jí)間連接如圖9.12所示,粗線和細(xì)線分別組成一個(gè)基本的連線圖形,因其形狀像蝴蝶,故稱為蝶形連接。分析圖9.12輸入/輸出節(jié)點(diǎn)編號(hào)間的關(guān)系可得出,將輸入節(jié)點(diǎn)的二進(jìn)制地址(編號(hào))的最高位與最低位互換位置,就獲得了與該輸入節(jié)點(diǎn)連接的輸出節(jié)點(diǎn)的二進(jìn)制地址,故其互連函數(shù)為圖9.12蝶形連接

例9.5多級(jí)立方體網(wǎng)也是動(dòng)態(tài)互連網(wǎng)絡(luò),如圖9.13所示。端口數(shù)為N,采用兩功能交換開(kāi)關(guān)(直通、交換),使用級(jí)挖或單元控制方式;互連模式為輸入級(jí)恒等置換、各級(jí)間子蝶形置換、輸出級(jí)逆混洗置換。試分析該網(wǎng)絡(luò)的互連函數(shù)。

解子蝶形置換互連函數(shù)為

圖9.13多級(jí)立方體網(wǎng)絡(luò)

MIN所連接的終端節(jié)點(diǎn)設(shè)備集中于網(wǎng)絡(luò)周圍,因此稱為集中式交換網(wǎng)絡(luò)。MIN通過(guò)一組開(kāi)關(guān)(交換開(kāi)關(guān))間接連接兩終端節(jié)點(diǎn)設(shè)備,因此MIN也稱為間接網(wǎng)絡(luò)。現(xiàn)在MIN已成為并行超級(jí)計(jì)算機(jī)、對(duì)稱多處理器系統(tǒng)、多計(jì)算機(jī)集群等高性能計(jì)算機(jī)中普遍采用的通信中樞。

9.3.4分布式交換網(wǎng)絡(luò)

用于分布式交換網(wǎng)絡(luò)的拓?fù)浣Y(jié)構(gòu)與集中式交換網(wǎng)絡(luò)有幾點(diǎn)不同:

①分布式交換開(kāi)關(guān)的數(shù)量與系統(tǒng)中的節(jié)點(diǎn)數(shù)量相同;

②所有在網(wǎng)絡(luò)中的節(jié)點(diǎn)用專用鏈路連接;

③是一種完全連接的拓?fù)浣Y(jié)構(gòu),可以提供最佳的連通性(即全連通);

④比交叉開(kāi)關(guān)網(wǎng)絡(luò)昂貴。大多數(shù)分布式交換網(wǎng)絡(luò)是靜態(tài)互連網(wǎng)絡(luò),因?yàn)樵垂?jié)點(diǎn)與目的節(jié)點(diǎn)之間的路徑是確定的,而靜態(tài)互連網(wǎng)絡(luò)一般用于可預(yù)知通信模式的、直接點(diǎn)對(duì)點(diǎn)的連接中。

1.線性網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

圖9.14是一組線性網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu),這組網(wǎng)絡(luò)的共同特點(diǎn)是從源節(jié)點(diǎn)到目的節(jié)點(diǎn)只有一條路徑。圖中只畫(huà)出了鏈路(邊)和交換節(jié)點(diǎn)(點(diǎn)),節(jié)點(diǎn)設(shè)備需要通過(guò)網(wǎng)絡(luò)接口連接到交換節(jié)點(diǎn)上。圖9.14線性網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

2.環(huán)形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

圖9.15是一組環(huán)形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu),與線性網(wǎng)絡(luò)結(jié)構(gòu)的最大不同是網(wǎng)絡(luò)中有一個(gè)連接所有節(jié)點(diǎn)的環(huán)。環(huán)形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)的最大特點(diǎn)是具有對(duì)稱性(Symmetry),即從任何節(jié)點(diǎn)看到的拓?fù)浣Y(jié)構(gòu)都是一樣的。對(duì)稱拓?fù)涞膬?yōu)點(diǎn)是網(wǎng)絡(luò)中所有節(jié)點(diǎn)具有相同的連接模式,每個(gè)節(jié)點(diǎn)或每條鏈路上的負(fù)載(或流量)是均勻分布的,網(wǎng)絡(luò)的硬件實(shí)現(xiàn)和軟件編程較容易。圖9.15環(huán)形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

3.網(wǎng)格形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

圖9.16是一組比較流行的網(wǎng)格形(Grid)網(wǎng)絡(luò)結(jié)構(gòu),屬于二維網(wǎng)絡(luò),有多種變體形式。網(wǎng)格結(jié)構(gòu)有規(guī)律,易于擴(kuò)展,其直徑與節(jié)點(diǎn)數(shù)的平方根成正比。許多大型商用系統(tǒng)都使用這種結(jié)構(gòu),如IlliacIV、MPP、DAP、CM-2和IntelParagon等。圖9.16網(wǎng)格形網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

4.超立方體網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

圖9.17是一組超立方體(Hypercube)網(wǎng)絡(luò),也稱為二進(jìn)制n-立方(Binaryn-Cube)網(wǎng)。n-立方網(wǎng)有N=2n個(gè)節(jié)點(diǎn),維數(shù)為n,節(jié)點(diǎn)分布在n維空間,節(jié)點(diǎn)度為n,鏈路數(shù)為N(lbN)/2,直徑為n,對(duì)分寬度為N/2。圖9.17超立方體網(wǎng)絡(luò)拓?fù)浣Y(jié)構(gòu)

9.4向量處理機(jī)

物理學(xué)和工程學(xué)的許多問(wèn)題都涉及陣列或其他高度規(guī)則的數(shù)據(jù)結(jié)構(gòu),經(jīng)常要在同一時(shí)間對(duì)不同數(shù)據(jù)集合完成相同運(yùn)算。數(shù)據(jù)的高度規(guī)則和程序的結(jié)構(gòu)化使通過(guò)并行執(zhí)行指令來(lái)加速程序執(zhí)行變得十分容易。采用向量處理機(jī)(VectorProcessors)可以有效地提高大型科學(xué)計(jì)算程序的執(zhí)行速度。

向量處理機(jī)在指令級(jí)并行計(jì)算機(jī)之前就已成功地商業(yè)化,它提供向量操作的高級(jí)運(yùn)算。向量處理機(jī)的典型特征是執(zhí)行向量指令,它可以解決指令級(jí)并行面臨的一些問(wèn)題。向量指令具有以下幾個(gè)重要特性:

(1)一條向量指令規(guī)定了相當(dāng)于一個(gè)完整循環(huán)所做的工作,每條指令代表數(shù)十或數(shù)百個(gè)操作,因此指令獲取和譯碼的帶寬必須足夠?qū)挕?/p>

(2)向量指令使結(jié)果的向量元素計(jì)算互不相關(guān),因此硬件在一條向量指令執(zhí)行期內(nèi)不必檢查數(shù)據(jù)相關(guān)(冒險(xiǎn))。可以利用并行的功能部件陣列、單一的深度流水的功能部件或者并行和流水的功能部件的組合來(lái)計(jì)算向量的各元素。

(3)硬件僅需要在兩條向量指令之間對(duì)每個(gè)向量操作數(shù)檢查一次數(shù)據(jù)相關(guān),而不需要對(duì)向量的每個(gè)元素進(jìn)行檢查。這意味著兩條向量指令之間需要的相關(guān)性檢查邏輯大致與兩條標(biāo)量指令之間需要的相同。

(4)如果向量元素是全部毗鄰的,則向量指令訪問(wèn)存儲(chǔ)器的最好模式是從一組交叉存取的存儲(chǔ)塊中獲取向量,這樣,對(duì)整個(gè)向量而言僅有一次等待主存的代價(jià)。

(5)由于整個(gè)循環(huán)由預(yù)先確定行為的向量指令替代,因此通常由循環(huán)分支產(chǎn)生的控制相關(guān)(冒險(xiǎn))是不存在的。

所以,向量運(yùn)算比具有相同數(shù)據(jù)項(xiàng)的標(biāo)量運(yùn)算更快。向量處理機(jī)主要用于大型科學(xué)和工程計(jì)算中。

9.4.1基本的向量體系結(jié)構(gòu)

向量處理器一般由一個(gè)普通的流水線標(biāo)量單元加上一個(gè)向量單元組成。向量單元內(nèi)的所有功能部件有幾個(gè)時(shí)鐘周期的等待時(shí)間(Latency),這樣,更短時(shí)鐘周期就可以與具有深度流水而不產(chǎn)生冒險(xiǎn)的較長(zhǎng)運(yùn)行時(shí)間的向量運(yùn)算相適應(yīng)。標(biāo)量單元基本上與先進(jìn)的流水線CPU相同,且商用向量機(jī)已內(nèi)置了亂序標(biāo)量單元(NECSX/5)和VLIW標(biāo)量單元(FujitsuVPP5000)。多數(shù)向量處理器允許以浮點(diǎn)數(shù)、整數(shù)或者邏輯數(shù)據(jù)處理向量。

向量-寄存器處理器的基本組成如圖9.18所示,它是一個(gè)以Cra量部分是MIPS,向量部分是MIPS的邏輯向量擴(kuò)展,其主要模塊功能y-如1下為:

圖9.18基本的向量寄存器體系結(jié)構(gòu)

(1)向量寄存器組。該結(jié)構(gòu)中有8個(gè)向量寄存器,每個(gè)向量寄存器是一個(gè)保存單一向量的定長(zhǎng)存儲(chǔ)塊,保存64個(gè)元素,并有至少兩個(gè)讀端口和一個(gè)寫(xiě)端口。

(2)向量功能單元。每個(gè)單元都是充分流水的,且在每個(gè)時(shí)鐘周期啟動(dòng)一次新操作。需要控制單元檢測(cè)來(lái)自功能部件沖突(結(jié)構(gòu)冒險(xiǎn))和來(lái)自寄存器訪問(wèn)沖突(數(shù)據(jù)冒險(xiǎn))產(chǎn)生的冒險(xiǎn)。

(3)向量Load-Store部件。這是將向量存入或取出存儲(chǔ)器的向量存儲(chǔ)控制單元。向量處理器的向量加載和存儲(chǔ)是完全流水的,在一個(gè)初始等待時(shí)間之后,數(shù)據(jù)可以在向量寄存器和存儲(chǔ)器之間以每個(gè)時(shí)鐘周期1個(gè)字的帶寬進(jìn)行傳送。這個(gè)部件通常也處理標(biāo)量加載和存儲(chǔ)。

(4)標(biāo)量寄存器組。標(biāo)量寄存器也能夠?yàn)橄蛄抗δ懿考峁┹斎霐?shù)據(jù),同時(shí)計(jì)算傳遞給向量Load-Store部件的地址。它們是常規(guī)32個(gè)通用寄存器和32個(gè)MIPS浮點(diǎn)寄存器。標(biāo)量值從標(biāo)量寄存器文件讀出,然后鎖存在向量功能部件的一個(gè)輸入端。圖9.19計(jì)算DAXPY算法

9.4.2現(xiàn)代的向量超級(jí)計(jì)算機(jī)CrayX1

CrayX1有與眾不同的處理器結(jié)構(gòu)。一個(gè)大的多流處理器(Multi-StreamingProcessor,MSP)由4個(gè)單流處理器(Single-StreamingProcessor,SSP)聯(lián)合組成,如圖9.20所示。每個(gè)SSP是一個(gè)完全的單片向量微處理器,包含一個(gè)標(biāo)量單元、數(shù)個(gè)標(biāo)量Cache和一個(gè)雙通道向量單元。SSP標(biāo)量單元是一個(gè)雙發(fā)射、亂序的超標(biāo)量處理器,有一個(gè)16KB的指令Cache和一個(gè)16KB的標(biāo)量寫(xiě)直達(dá)(Write-Through)數(shù)據(jù)Cache,兩者的兩路設(shè)置與32字節(jié)的Cache塊相關(guān)聯(lián)。

SSP向量單元包含1個(gè)向量寄存器文件、3個(gè)向量運(yùn)算單元和1個(gè)向量Load-Store控制部件。向量功能單元深度流水比超標(biāo)量發(fā)射機(jī)制更容易實(shí)現(xiàn),因此CraX1向量單元以標(biāo)量單元(400MHz)兩倍的時(shí)鐘頻率(800MHz)運(yùn)行。每條通道每個(gè)時(shí)鐘周y期能夠完成一個(gè)64位浮點(diǎn)加法和一個(gè)64位浮點(diǎn)乘法,使每個(gè)MSP達(dá)到12.8GFLOPS的峰值性能。

圖9.20CrayMSP模塊

在CrayX1封裝層次的上一級(jí)(如圖9.21所示),4個(gè)MSP、16個(gè)存儲(chǔ)控制器芯片和DRAM一起被放置在一塊印制電路板上,形成一個(gè)X1節(jié)點(diǎn)(Node)。每個(gè)存儲(chǔ)控制器芯片有8個(gè)獨(dú)立的RambusDRAM信道,每條信道提供1.6GB/s的存儲(chǔ)器帶寬。交叉全部的128條存儲(chǔ)器信道,節(jié)點(diǎn)有超過(guò)200GB/s的主存帶寬。圖9.21CrayX1節(jié)點(diǎn)

9.5多處理器系統(tǒng)

9.5.1UMA對(duì)稱多處理器系統(tǒng)UMA對(duì)稱多處理器系統(tǒng)是迄今為止最流行的并行系統(tǒng)結(jié)構(gòu),它采用集中式共享存儲(chǔ)器結(jié)構(gòu)。由于共享存儲(chǔ)器對(duì)每個(gè)處理器而言都是對(duì)等的,并且每個(gè)處理器訪問(wèn)存儲(chǔ)器的時(shí)間相同,所以,UMA對(duì)稱多處理器系統(tǒng)具有存儲(chǔ)器對(duì)稱、訪問(wèn)時(shí)間一致的特點(diǎn)。

UMA對(duì)稱多處理器系統(tǒng)的體系結(jié)構(gòu)如圖9.22所示,多個(gè)處理器Cache子系統(tǒng)共享同一個(gè)物理存儲(chǔ)器,存儲(chǔ)器可以按多組方式組織。

圖9.22集中式共享存儲(chǔ)器多處理器基本結(jié)構(gòu)

9.5.2多處理器系統(tǒng)的Cache一致性

假設(shè)有兩個(gè)處理器CPU1和CPU2分別擁有各自的寫(xiě)直達(dá)Cache,用Cache1和Cache2表示。當(dāng)CPU1和CPU2對(duì)同一個(gè)存儲(chǔ)器單元X進(jìn)行讀/寫(xiě)操作時(shí),如果X的值被CPU1改寫(xiě),那么Cache1和存儲(chǔ)器中的副本會(huì)被更新,但Cache2卻未更新,這樣,在Cache1和Cache2中就有X單元的兩個(gè)不同的副本,出現(xiàn)數(shù)據(jù)在Cache中不一致的狀況。如果此時(shí)CPU2讀取X,則它將從Cache2中得到X單元的舊數(shù)據(jù),而不是X單元當(dāng)前最新的數(shù)據(jù),從而可能引起處理異常。這就是多處理器系統(tǒng)的Cache一致性(CacheCoherence或CacheConsistency)問(wèn)題。

在具有多級(jí)Cache的單處理器系統(tǒng)和具有共享存儲(chǔ)器的多處理器系統(tǒng)中都存在Cache一致性的問(wèn)題,特別是處理器眾多時(shí),Cache一致性問(wèn)題還是一個(gè)嚴(yán)重而棘手的問(wèn)題。產(chǎn)生Cache不一致性的原因主要有以下三種:

(1)共享數(shù)據(jù)。為了減少訪問(wèn)遠(yuǎn)程共享數(shù)據(jù)的時(shí)延(可通過(guò)共享數(shù)據(jù)遷移實(shí)現(xiàn))、減少對(duì)共享存儲(chǔ)器帶寬的需求和訪問(wèn)共享數(shù)據(jù)時(shí)的競(jìng)爭(zhēng)(可通過(guò)共享數(shù)據(jù)復(fù)制實(shí)現(xiàn)),在多個(gè)處理器上運(yùn)行的程序會(huì)要求在多個(gè)Cache中有同一個(gè)數(shù)據(jù)的副本。而這種共享性要求正是Cache不一致性的主要原因。

(2)進(jìn)程遷移。在多處理器系統(tǒng)中,進(jìn)程可以在處理器中相互遷移。如果某個(gè)處理器中的進(jìn)程修改了私有Cache中的數(shù)據(jù),但還沒(méi)寫(xiě)回主存前,由于某種原因需要遷移到其他處理器中繼續(xù)運(yùn)行,此時(shí)讀到的存儲(chǔ)器中的數(shù)據(jù)將是過(guò)時(shí)的數(shù)據(jù)。

(3)I/O操作。繞過(guò)Cache的I/O操作會(huì)引起Cache與共享主存的不一致性。例如,當(dāng)DMA控制器直接對(duì)主存進(jìn)行寫(xiě)操作時(shí),若Cache中有相應(yīng)數(shù)據(jù)的副本,就會(huì)造成主存與Cache之間的不一致。

為了維護(hù)Cache一致性,多處理器系統(tǒng)引入了Cache一致性協(xié)議(CacheCoherenceProtocol),該協(xié)議是由Cache、CPU和共享存儲(chǔ)器共同實(shí)現(xiàn)的防止多個(gè)Cache中出現(xiàn)同一數(shù)據(jù)不同副本的規(guī)則集合而構(gòu)成的。協(xié)議有多種,但目的只有一個(gè):防止在兩個(gè)或者更多的Cache中出現(xiàn)同一塊數(shù)據(jù)的不同版本?？紤]到快速性,多處理器系統(tǒng)一般通過(guò)硬件實(shí)現(xiàn)Cache一致性協(xié)議。

實(shí)現(xiàn)Cache一致性協(xié)議的關(guān)鍵在于跟蹤所有共享數(shù)據(jù)塊的狀態(tài)。目前廣泛采用的有兩類協(xié)議,它們采用不同的技術(shù)跟蹤共享數(shù)據(jù)。

(1)監(jiān)聽(tīng)式:主要用于總線作為互連網(wǎng)絡(luò)的系統(tǒng)。該方案中,Cache控制器通過(guò)監(jiān)聽(tīng)總線行為來(lái)決定自己將采取哪種行動(dòng)。這種Cache被稱為監(jiān)聽(tīng)型Cache(SnoopyCache),該方案可用于具有廣播特性的通信媒介的系統(tǒng)中。

(2)目錄式:該方案是把共享存儲(chǔ)器中共享數(shù)據(jù)塊的狀態(tài)及相關(guān)信息存放在一個(gè)目錄(Directory)中,通過(guò)訪問(wèn)目錄來(lái)跟蹤所有共享數(shù)據(jù)塊的狀態(tài)。目錄式一致性協(xié)議比監(jiān)聽(tīng)式的實(shí)現(xiàn)開(kāi)銷略大些,但是目錄式協(xié)議可以用來(lái)擴(kuò)展更多的處理器。Sun公司的T1處理器設(shè)計(jì)就采用了目錄式協(xié)議。

1.監(jiān)聽(tīng)協(xié)議(SnoopingProtocol)

大多數(shù)多處理器系統(tǒng)使用微處理器,并采用總線將微處理器和它的Cache與單一的共享存儲(chǔ)器進(jìn)行連接,這使得監(jiān)聽(tīng)協(xié)議可以使用已有的物理連接總線來(lái)查詢Cache塊的狀態(tài)。

有兩種常用的監(jiān)聽(tīng)式Cache一致性協(xié)議:寫(xiě)直達(dá)協(xié)議和MESICache一致性協(xié)議。

1)寫(xiě)直達(dá)協(xié)議

寫(xiě)直達(dá)協(xié)議是最簡(jiǎn)單的Cache一致性協(xié)議,該協(xié)議保證共享存儲(chǔ)器的數(shù)據(jù)總是最新的。該協(xié)議規(guī)定:

(1)當(dāng)CPU要讀的字不在Cache中時(shí),發(fā)生讀缺失,Cache控制器則把包括該字的一個(gè)數(shù)據(jù)塊從共享存儲(chǔ)器讀入Cache(Cache塊一般為32或64字節(jié))。接下來(lái)對(duì)該數(shù)據(jù)塊的讀操作就直接在Cache中進(jìn)行(即讀命中)。

(2)當(dāng)發(fā)生寫(xiě)缺失時(shí),被修改的字寫(xiě)入共享存儲(chǔ)器,但包括該字的數(shù)據(jù)塊不調(diào)入Cache。當(dāng)發(fā)生寫(xiě)命中時(shí),修改Cache的同時(shí)把該字直接寫(xiě)入共享存儲(chǔ)器。

2)MESICache一致性協(xié)議

為了保證一定的總線流量,人們?cè)O(shè)計(jì)了MESI協(xié)議(也稱為寫(xiě)回協(xié)議),即:

(1)CPU做寫(xiě)操作時(shí),修改數(shù)據(jù)寫(xiě)入Cache塊,但不立刻寫(xiě)入共享存儲(chǔ)器,在Cache中設(shè)置狀態(tài)以表示該Cache塊中的數(shù)據(jù)是正確的,共享存儲(chǔ)器中的數(shù)據(jù)是過(guò)時(shí)的。

(2)僅當(dāng)需要替換該Cache塊時(shí),將該Cache塊寫(xiě)回共享存儲(chǔ)器。

MESI協(xié)議規(guī)定每個(gè)Cache塊都處于以下四種狀態(tài)之一:

·無(wú)效(Invalid)———該Cache塊包含的數(shù)據(jù)無(wú)效。

·共享(Shared)———多個(gè)Cache中都有這塊數(shù)據(jù),共享存儲(chǔ)器中的數(shù)據(jù)是最新的。

·獨(dú)占(Exclusive)———沒(méi)有其他Cache包括這塊數(shù)據(jù),共享存儲(chǔ)器中的數(shù)據(jù)是最新的。

·修改(Modified)———該塊數(shù)據(jù)是有效的,共享存儲(chǔ)器中的數(shù)據(jù)是無(wú)效的,而且在其他Cache中沒(méi)有該數(shù)據(jù)塊的副本。

(1)“無(wú)效”狀態(tài)的設(shè)置。當(dāng)系統(tǒng)啟動(dòng)時(shí),系統(tǒng)中所有的CPU通過(guò)它的Cache控制器將所有Cache塊標(biāo)記為I(無(wú)效),并啟動(dòng)Cache控制器開(kāi)始監(jiān)聽(tīng)總線。

(2)“獨(dú)占”狀態(tài)的設(shè)置。假設(shè)CPU1首先讀取共享存儲(chǔ)器中的數(shù)據(jù)塊D,此時(shí)Cache1控制器不能監(jiān)聽(tīng)到其他Cache擁有數(shù)據(jù)塊D的消息,因?yàn)檫@是當(dāng)前Cache中唯一的一個(gè)副本,所以Cache1控制器將數(shù)據(jù)塊D讀入Cache1后,將其所在的Cache塊標(biāo)記為E(獨(dú)占)。之后,CPU1對(duì)在該Cache塊中的數(shù)據(jù)進(jìn)行訪問(wèn)時(shí),均不需要經(jīng)過(guò)總線。

(3)“共享”狀態(tài)的設(shè)置。在Cache1獨(dú)占數(shù)據(jù)塊D期間,如果CPU2需要讀取數(shù)據(jù)塊D,那么它要先向總線發(fā)送讀數(shù)據(jù)塊D的消息(在總線上提供數(shù)據(jù)塊D的地址),并啟動(dòng)Cache2控制器將數(shù)據(jù)塊D從共享存儲(chǔ)器讀入到Cache2,同時(shí)監(jiān)聽(tīng)總線。數(shù)據(jù)塊D的其他持有者(如Cache1)監(jiān)聽(tīng)到CPU2讀數(shù)據(jù)塊D的消息后,立刻在總線上發(fā)布通告,宣稱自己擁有一份該數(shù)據(jù)的副本,并將數(shù)據(jù)塊D所在的Cache塊標(biāo)記為S(共享)。同時(shí),Cache2控制器監(jiān)聽(tīng)到該總線通告后,將Cache2中數(shù)據(jù)塊D所在的Cache塊也標(biāo)記為S。這樣,數(shù)據(jù)塊D在Cache中的所有副本就都被標(biāo)記成S狀態(tài)。S狀態(tài)意味著該數(shù)據(jù)塊在多個(gè)Cache中存在,并且共享存儲(chǔ)器中的數(shù)據(jù)是最新的。CPU對(duì)處于S狀態(tài)的Cache塊進(jìn)行讀操作時(shí)不使用總線。

(4)“修改”狀態(tài)的設(shè)置。如果CPU2需要向S狀態(tài)的Cache塊寫(xiě)入數(shù)據(jù),它會(huì)把一個(gè)無(wú)效信號(hào)和寫(xiě)入字的地址通過(guò)總線傳送(廣播)給其他的CPU,通知它們把相應(yīng)的數(shù)據(jù)副本置為I(無(wú)效)。而CPU2自己的Cache塊的狀態(tài)則改變成M(修改),但該塊并不需要寫(xiě)回共享存儲(chǔ)器。如果處于E狀態(tài)的Cache塊發(fā)生了寫(xiě)操作,則不需要給其他的Cache發(fā)送無(wú)效信號(hào),因?yàn)樗荂ache中唯一的副本。

(5)狀態(tài)轉(zhuǎn)換。如果數(shù)據(jù)塊D已在Cache2中,且它所在的Cache塊處于M狀態(tài),此時(shí),當(dāng)CPU3讀數(shù)據(jù)塊D時(shí),擁有該塊的CPU2知道共享存儲(chǔ)器中的數(shù)據(jù)已是無(wú)效,因此CPU2就向總線發(fā)送將該塊寫(xiě)回共享存儲(chǔ)器的消息。CPU3監(jiān)聽(tīng)到CPU2的寫(xiě)消息后,就會(huì)等待CPU2把該塊寫(xiě)回共享存儲(chǔ)器后,CPU3再?gòu)墓蚕泶鎯?chǔ)器中取得數(shù)據(jù)塊D的副本,然后Cache2和Cache3中數(shù)據(jù)塊D所在的Cache塊都被標(biāo)記為S。

如果數(shù)據(jù)塊D已在Cache2中,且它所在的Cache塊處于M狀態(tài),此時(shí),當(dāng)CPU1想向數(shù)據(jù)塊D中寫(xiě)入一個(gè)字時(shí),CPU1向總線發(fā)出寫(xiě)請(qǐng)求。CPU2監(jiān)聽(tīng)到該寫(xiě)請(qǐng)求后,也向總線發(fā)送一個(gè)消息,通知CPU1等待它把該塊寫(xiě)回共享存儲(chǔ)器。當(dāng)寫(xiě)回操作完成后,CPU2將Cache2中的該塊標(biāo)記為I,因?yàn)樗繡PU1將會(huì)修改該塊。如果使用寫(xiě)分配策略,CPU1僅修改Cache1中相應(yīng)塊,并標(biāo)記為M狀態(tài)。如果沒(méi)有使用寫(xiě)分配策略,將直接對(duì)共享存儲(chǔ)器執(zhí)行寫(xiě)操作而該塊不會(huì)被讀入任何Cache。圖9.23為MESI協(xié)議的簡(jiǎn)化狀態(tài)轉(zhuǎn)換圖。圖9.23MESI協(xié)議簡(jiǎn)化狀態(tài)轉(zhuǎn)換圖

還有一種不常使用的MOESI擴(kuò)展協(xié)議。除了上述的4種獨(dú)立的狀態(tài)之外,MOESI還增加了“擁有者”(Owner)狀態(tài),見(jiàn)表9.1。在MOESI協(xié)議中,如果CPU1獲得了共享存儲(chǔ)器某個(gè)數(shù)據(jù)區(qū)的“擁有者”狀態(tài),那么其他處理器可以直接從CPU1的緩存中進(jìn)行更新,而無(wú)須訪問(wèn)共享存儲(chǔ)區(qū)。這個(gè)過(guò)程減少了對(duì)共享存儲(chǔ)區(qū)的訪問(wèn)次數(shù),加快了訪問(wèn)的速度。在Futurebus+總線中研發(fā)的MOESI協(xié)議也用在了AMD的多處理器芯片組AMD760MP中。

2.目錄協(xié)議(DirectoryProtocol)

目錄協(xié)議不需要向所有Cache進(jìn)行廣播,因而便于處理器速度和數(shù)量的增加,在大規(guī)模系統(tǒng)中越來(lái)越受到青睞。目錄協(xié)議經(jīng)常用在分布式存儲(chǔ)器結(jié)構(gòu)的系統(tǒng)中,但也適用于按組進(jìn)行組織的集中式存儲(chǔ)器結(jié)構(gòu)系統(tǒng)。

目錄是一個(gè)在專用存儲(chǔ)器中存儲(chǔ)的數(shù)據(jù)結(jié)構(gòu),它記錄著調(diào)入Cache的每個(gè)數(shù)據(jù)塊的訪問(wèn)狀態(tài)、該塊在各個(gè)處理器的共享狀態(tài)以及是否修改過(guò)等信息。

目錄協(xié)議的設(shè)計(jì)思想是利用目錄來(lái)維護(hù)Cache塊的共享信息。最簡(jiǎn)單的目錄協(xié)議實(shí)現(xiàn)機(jī)制是在目錄中給每個(gè)共享存儲(chǔ)器數(shù)據(jù)塊分配一個(gè)目錄項(xiàng),目錄項(xiàng)個(gè)數(shù)由存儲(chǔ)器中塊的個(gè)數(shù)(其中每個(gè)塊的大小和二級(jí)或三級(jí)的Cache塊大小一樣)和處理器個(gè)數(shù)的乘積決定。由于每條訪問(wèn)共享存儲(chǔ)器的指令都要查詢這個(gè)目錄,因此這個(gè)目錄必須保存在速度非?？斓膶Ｓ糜布?以保證在不到一個(gè)總線周期的時(shí)間內(nèi)做出響應(yīng)。

為了達(dá)到Cache一致性要求,目錄協(xié)議需要完成兩個(gè)基本操作:①處理讀缺失;②處理共享、未修改的Cache塊的寫(xiě)操作(處理共享數(shù)據(jù)塊的寫(xiě)缺失由這兩個(gè)操作簡(jiǎn)單組合而成)。為了實(shí)現(xiàn)這些操作,目錄必須跟蹤每個(gè)Cache塊的狀態(tài)。最簡(jiǎn)單的Cache塊狀態(tài)可定義為以下三種:

(1)共享:在一個(gè)或多個(gè)處理器上具有這個(gè)數(shù)據(jù)塊的副本,且主存中的值是最新值(所有Cache副本均相同)。

(2)未緩存:所有處理器的Cache中都沒(méi)有該數(shù)據(jù)塊的副本。

(3)獨(dú)占:僅有一個(gè)處理器上有該數(shù)據(jù)塊的副本,且已對(duì)該塊進(jìn)行了寫(xiě)操作,而主存的副本仍是舊的(無(wú)效)。這個(gè)處理器稱為該塊的擁有者。

圖9.24是包含位向量的基本目錄結(jié)構(gòu),每個(gè)目錄項(xiàng)對(duì)應(yīng)一個(gè)存儲(chǔ)塊,且由狀態(tài)和位向量N組成。狀態(tài)描述該目錄項(xiàng)所對(duì)應(yīng)存儲(chǔ)塊的當(dāng)前情況(至少用上述3種狀態(tài)表示);位向量共有=2n位,每一位對(duì)應(yīng)一個(gè)處理器的本地Cache,用于指出該Cache中有無(wú)該存儲(chǔ)塊的副本。圖9.24分布在節(jié)點(diǎn)上的目錄結(jié)構(gòu)

在基于目錄的協(xié)議中,目錄利用消息機(jī)制來(lái)完成一致性協(xié)議所要求的操作功能。當(dāng)CPU對(duì)共享存儲(chǔ)器操作出現(xiàn)讀缺失、寫(xiě)缺失、數(shù)據(jù)寫(xiě)回狀況時(shí),由MMU(MemoryManagementUnit,存儲(chǔ)器管理單元)向目錄發(fā)送一個(gè)消息,目錄硬件按照消息的請(qǐng)求進(jìn)行服務(wù),并更新目錄狀態(tài)。具體實(shí)現(xiàn)如下:

(1)當(dāng)一個(gè)數(shù)據(jù)塊處于未緩存狀態(tài)時(shí),若目錄接收到對(duì)此塊讀操作而引發(fā)的讀缺失請(qǐng)求,則目錄硬件會(huì)將存儲(chǔ)器數(shù)據(jù)送往請(qǐng)求方處理器,且本處理器成為此塊的唯一共享節(jié)點(diǎn),本塊的狀態(tài)轉(zhuǎn)換為共享。

(2)當(dāng)一個(gè)數(shù)據(jù)塊是共享狀態(tài),且存儲(chǔ)器中的數(shù)據(jù)是當(dāng)前最新值時(shí),若目錄接收到對(duì)此塊讀操作而引發(fā)的讀缺失請(qǐng)求,則目錄硬件會(huì)將存儲(chǔ)器數(shù)據(jù)送往請(qǐng)求方處理器,并將其加入共享集合(如位向量)。

(3)當(dāng)某數(shù)據(jù)塊處于獨(dú)占狀態(tài)時(shí),本塊的最新值保存在共享集合指出的擁有者處理器中,這時(shí)有三種可能的目錄請(qǐng)求消息。若為讀缺失消息,則目錄硬件會(huì)將“取數(shù)據(jù)”的消息發(fā)往擁有者處理器,使該數(shù)據(jù)塊的狀態(tài)轉(zhuǎn)變?yōu)楣蚕?并將所獲得數(shù)據(jù)通過(guò)目錄寫(xiě)入存儲(chǔ)器,再把該數(shù)據(jù)返送請(qǐng)求方處理器,將請(qǐng)求方處理器加入共享集合。

9.5.3NUMA對(duì)稱多處理器系統(tǒng)

對(duì)于處理器數(shù)目較少的多處理器系統(tǒng),各個(gè)處理器可以共享單個(gè)集中式存儲(chǔ)器。在使用大容量Cache的情況下,單一存儲(chǔ)器(可能是多組)能夠確保小數(shù)目處理器的存儲(chǔ)訪問(wèn)得到及時(shí)響應(yīng)。通過(guò)使用多個(gè)點(diǎn)對(duì)點(diǎn)連接,或者通過(guò)交換機(jī),集中共享存儲(chǔ)器(組)設(shè)計(jì)可以擴(kuò)展到幾十個(gè)處理器。隨著處理器數(shù)量的增多以及處理器對(duì)存儲(chǔ)器要求的增加,系統(tǒng)的任何集中式資源都會(huì)變成瓶頸。

為了支持更多的處理器,減少存儲(chǔ)器在為多個(gè)處理器提供所需要的帶寬時(shí)無(wú)法避免的較長(zhǎng)時(shí)延,存儲(chǔ)器不能再按照集中共享方式組織,而必須分布于各個(gè)處理器中,圖9.25就是基于分布式存儲(chǔ)器(DM)結(jié)構(gòu)的多處理器系統(tǒng)的組織結(jié)構(gòu)。將存儲(chǔ)器分布到各個(gè)節(jié)點(diǎn)上有兩個(gè)好處:①如果大部分訪問(wèn)是在節(jié)點(diǎn)內(nèi)的本地存儲(chǔ)器中進(jìn)行的,則這樣做是增大存儲(chǔ)器帶寬比較經(jīng)濟(jì)的方法;②縮短了本地存儲(chǔ)器訪問(wèn)的時(shí)延。分布式存儲(chǔ)器系統(tǒng)結(jié)構(gòu)的主要缺點(diǎn)是:由于處理器不再共享單一集中存儲(chǔ)器,處理器間的數(shù)據(jù)通信在某種程度上變得更加復(fù)雜,且時(shí)延也更大。圖9.25DM-MIMD系統(tǒng)的基本結(jié)構(gòu)

在分布式存儲(chǔ)器結(jié)構(gòu)的系統(tǒng)中,存儲(chǔ)器有兩種組織結(jié)構(gòu)。第一種是所有分布的存儲(chǔ)器共享地址空間。物理上分開(kāi)的存儲(chǔ)器能夠作為邏輯上共享的地址空間進(jìn)行尋址,就是說(shuō)只要有正確的訪問(wèn)權(quán)限,任何一個(gè)處理器都能夠通過(guò)引用地址的方式訪問(wèn)任意節(jié)點(diǎn)上的存儲(chǔ)器。這類機(jī)器稱為分布式共享存儲(chǔ)器(DSM)系統(tǒng)。所謂共享存儲(chǔ)器,指的是共享尋址空間,即兩個(gè)處理器中相同的物理地址指向存儲(chǔ)器中的同一個(gè)存儲(chǔ)位置。與對(duì)稱式共享存儲(chǔ)器多處理器系統(tǒng)(即UMA)相比,DSM系統(tǒng)由于訪問(wèn)時(shí)間取決于數(shù)據(jù)在存儲(chǔ)器中的位置,因而也稱為NUMA(非一致性存儲(chǔ)器訪問(wèn))。

第二種是存儲(chǔ)器地址空間由多個(gè)獨(dú)立私有的地址空間組成,這些私有地址空間在邏輯上是分散的,并且不能被遠(yuǎn)程處理器直接尋址。在這種機(jī)器中,兩個(gè)不同處理器中相同的物理地址分別指向兩個(gè)不同的存儲(chǔ)器,每個(gè)處理器-存儲(chǔ)器模塊本質(zhì)上是一臺(tái)獨(dú)立的計(jì)算機(jī)。采用這種存儲(chǔ)器結(jié)構(gòu)的計(jì)算機(jī)起初由不同的處理節(jié)點(diǎn)和專用互連網(wǎng)絡(luò)組成,而現(xiàn)在已演變?yōu)榧骸?/p>

每一種地址空間組織方式都有相應(yīng)的通信機(jī)制。對(duì)于共享地址空間的計(jì)算機(jī)系統(tǒng),通過(guò)Load和Store操作隱式地在處理器間傳遞數(shù)據(jù)。對(duì)于有多個(gè)尋址空間的計(jì)算機(jī)系統(tǒng),通過(guò)顯式地在處理器間傳送消息來(lái)完成數(shù)據(jù)通信,為此,這類機(jī)器也稱為消息傳遞計(jì)算機(jī)系統(tǒng)。

NUMA系統(tǒng)與其他多處理器系統(tǒng)相比有以下特點(diǎn):

(1)所有CPU看到的是同一個(gè)地址空間。

(2)使用Load和Store指令訪問(wèn)遠(yuǎn)程共享存儲(chǔ)器。

(3)訪問(wèn)遠(yuǎn)程共享存儲(chǔ)器比訪問(wèn)本地共享存儲(chǔ)器速度慢。

NUMA系統(tǒng)的一個(gè)重要特性是NUMA因子(NUMAFactor),NUMA因子表示從本地與非本地存儲(chǔ)單元存取數(shù)據(jù)的延遲差。依賴于系統(tǒng)的連接結(jié)構(gòu),系統(tǒng)不同部分的NUMA因子可以不同:從相鄰節(jié)點(diǎn)存取數(shù)據(jù)要比從相對(duì)較遠(yuǎn)的節(jié)點(diǎn)存取數(shù)據(jù)快,因?yàn)檩^遠(yuǎn)的節(jié)點(diǎn)可能要通過(guò)多級(jí)交換開(kāi)關(guān)存取數(shù)據(jù)。所以,當(dāng)提及NUMA因子時(shí),主要是指最大的網(wǎng)絡(luò)截面,即兩個(gè)處理器間的最大距離。

在圖9.26中,節(jié)點(diǎn)上的CPU由交叉開(kāi)關(guān)連接在存儲(chǔ)器的公共部分構(gòu)成SMP節(jié)點(diǎn),所有處理器都可以訪問(wèn)所有需要的地址空間,而使SMP節(jié)點(diǎn)共享存儲(chǔ)器的重要途徑就是使用COMA(Cache-OnlyMemoryArchitecture)或CC-NUMA(CacheCoherentNon-UniformMemoryAccess)結(jié)構(gòu)。因此,圖9.26描述的系統(tǒng)可以看成共享存儲(chǔ)器MIMD(SM-MIMD)系統(tǒng)。圖9.26SM-MIMD系統(tǒng)框圖

多處理器系統(tǒng)結(jié)構(gòu)涉及一個(gè)大而復(fù)雜的領(lǐng)域,其中很多領(lǐng)域仍處于不斷發(fā)展中,新的想法層出不窮,而失敗的系統(tǒng)結(jié)構(gòu)也屢見(jiàn)不鮮。多處理器系統(tǒng)的應(yīng)用非常廣泛,不論是運(yùn)行相互之間沒(méi)有通信的獨(dú)立任務(wù),還是必須通過(guò)線程通信才能完成的并行程序,都可以使用多處理器系統(tǒng)。

9.6圖形處理單元體系結(jié)構(gòu)

9.6.1簡(jiǎn)介隨著圖形顯示器的大規(guī)模應(yīng)用以及計(jì)算機(jī)游戲行業(yè)(包括PC和專用游戲機(jī))的快速增長(zhǎng),傳統(tǒng)的微處理器對(duì)于圖形處理力不從心,摩爾定律又限制了微處理器可用的晶體管數(shù)量,因此許多公司加大對(duì)圖形處理器的投資,也使得圖形處理器的發(fā)展速度比主流微處理器的發(fā)展速度更快。

圖9.27給出了已有CPU-GPU異構(gòu)系統(tǒng)的架構(gòu)。圖9.27(a)代表了分立GPU系統(tǒng)的典型布局。圖9.27(b)為AMDAthlon64架構(gòu)和隨后的IntelNehalem架構(gòu),通過(guò)將訪存控制器集成到CPU芯片上減少訪存延遲。圖9.27(c)為AMD加速處理單元(AcceleratedProcessorUint,APU),它將CPU和GPU集成到同一芯片上,這是一個(gè)巨大的突破,特別是隨著異構(gòu)統(tǒng)一內(nèi)存訪問(wèn)技術(shù)的引入,不但統(tǒng)一了CPU和GPU的內(nèi)存空間,而且還在兩者間保證了緩存的一致性。圖9.27CPU-GPU架構(gòu)

GPU僅僅是并行計(jì)算領(lǐng)域內(nèi)一個(gè)小的分支,但與CPU相比,GPU有以下優(yōu)點(diǎn):

(1)浮點(diǎn)計(jì)算能力強(qiáng)。

(2)超高性價(jià)比。

(3)綠色功耗比。

(4)普及度廣。

9.6.2GPU體系結(jié)構(gòu)

為了更好地認(rèn)知GPU體系結(jié)構(gòu),首先介紹幾個(gè)關(guān)于GPU體系結(jié)構(gòu)的相關(guān)術(shù)語(yǔ):

(1)流多處理器(StreamingMultiprocessors,SM):是GPU計(jì)算性能的基本單元之一。SM中包含SP、DP、SFU等,SM采用單指令多線程(SIMT)的執(zhí)行方式,實(shí)現(xiàn)向量化,保證多個(gè)線程可以同時(shí)執(zhí)行。

(2)流處理器(StreamingProcessor,SP):也稱Core,是GPU運(yùn)算的基本單元。早期GPU的流處理器僅支持單精度浮點(diǎn)運(yùn)算和整數(shù)運(yùn)算,隨著GPU體系結(jié)構(gòu)的不斷發(fā)展,使得GPU可以支持雙精度浮點(diǎn)運(yùn)算和完整的32位整數(shù)運(yùn)算。

(3)雙精度浮點(diǎn)運(yùn)算單元(DP):專用于雙精度浮點(diǎn)運(yùn)算的處理單元。

(4)特殊功能單元(SpecialFunctionUnit,SFU):用于執(zhí)行超越函數(shù)指令,比如正弦、余弦、倒數(shù)和平方根等函數(shù)。每個(gè)SFU一次執(zhí)行一個(gè)線程的一條指令需要一個(gè)時(shí)鐘周期,SFU并不始終占用SM中的分發(fā)單元,即當(dāng)SFU處于執(zhí)行狀態(tài)時(shí),指令分發(fā)單元可以向其他的執(zhí)行單元分發(fā)相應(yīng)的指令。

(5)線程處理器簇(ThreadProcessingCluster,TPC):由SM控制器、SM和L1Cache組成,存在于Tesla架構(gòu)和Pascal架構(gòu)中。G80架構(gòu)包含2個(gè)SM和16KBL1Cache,GT200架構(gòu)包含3個(gè)SM和24KBL1Cache。

(6)圖像處理器簇(GraphProcessingCluster,GPC):類似于TPC,是介于整個(gè)GPU和SM間的硬件單元。GPC由1個(gè)光柵單元、4個(gè)SM和4個(gè)SM控制器組成,且GPC中的SM數(shù)量是可擴(kuò)展的。

(7)流處理器陣列(Scalable/StreamingProcessorArray,SPA):所有處理核心和高速緩存的總和,包含所有的SM、TPC和GPC,與存儲(chǔ)器系統(tǒng)共同組成GPU架構(gòu)。

(8)存儲(chǔ)控制器(MemoryController,MMC):控制存儲(chǔ)訪問(wèn)的單元合并訪存,每個(gè)存儲(chǔ)控制器可以支持一定位寬的數(shù)據(jù)合并訪存。

(9)光柵控制單元(RasterOperationProcessors,ROP):每一個(gè)光柵控制單元都和特定的內(nèi)存分區(qū)配對(duì),對(duì)于游戲來(lái)說(shuō)主要負(fù)責(zé)計(jì)算機(jī)游戲的光線和反射運(yùn)算,兼顧全屏抗鋸齒、高分辨率、煙霧、火焰等效果。游戲中的光影效果越厲害,對(duì)光柵單元的性能要求也就越高。

(10)存取單元(Load/StoreUnits,LD/ST):負(fù)責(zé)全局內(nèi)存、共享內(nèi)存、常量?jī)?nèi)存、紋理加載等存儲(chǔ)器訪問(wèn)操作。

圖9.28所示為Pascal架構(gòu)GPU的SM示意圖。Pascal架構(gòu)(GP100)中的SM可以分為兩個(gè)處理塊,每個(gè)處理塊擁有32個(gè)單精度CUDA核、16個(gè)雙精度處理單元、8個(gè)LD/ST單元、8個(gè)SFU單元、1個(gè)指令緩沖器、1個(gè)Warp調(diào)度器、2個(gè)指令分發(fā)單元、32768個(gè)32位寄存器。整個(gè)SM還擁有4個(gè)紋理單元和64KB共享存儲(chǔ)器。圖9.28Pascal架構(gòu)SM示意圖

在Pascal架構(gòu)中,2個(gè)SM組成1個(gè)TPC,5個(gè)TPC組成1個(gè)GPC。一個(gè)完整的GP100(如圖9.29所示)核心擁有6個(gè)GPC、30個(gè)TPC、60個(gè)Pascal架構(gòu)的SM、8個(gè)512位寬的存儲(chǔ)控制器(共4096位)、3840個(gè)單精度CUDA核、240個(gè)紋理單元和4個(gè)NVLink單元。每個(gè)存儲(chǔ)控制器管理512KB的L2Cache(共4096KB)。兩個(gè)存儲(chǔ)控制器控制1個(gè)HBM2(HiBandwidthMemory2)DRAMStack。每個(gè)NVLink單元可以提供40GB/s的雙向通信帶寬。圖9.29完整的GP100GPU

NVLink可以將多個(gè)GPU交叉連接構(gòu)成網(wǎng)絡(luò),以獲得更好的通信效益。圖9.30展示了由8個(gè)P100GPU組成的立體網(wǎng)絡(luò)。圖中每條虛線表示的NVLink連接可以提供40GB/s的雙向通信帶寬。若是相互連接的GPU數(shù)量下降,那么多出來(lái)的NVLink單元可以與現(xiàn)有其他GPU連接,若是兩個(gè)GPU間連接了兩個(gè)NVLink,那么其雙向通信帶寬為80GB/s。圖9.308個(gè)P100組成的立體網(wǎng)格

相比Kepler架構(gòu)和Maxwell架構(gòu),Pascal架構(gòu)引入了不少新的技術(shù),提供了更加強(qiáng)大的性能。Pascal架構(gòu)的優(yōu)勢(shì)如下:

(1)浮點(diǎn)運(yùn)算性能強(qiáng)。TeslaP100GPU提供了5.3TFlops雙精度運(yùn)算性能、10.6TFlops和21.2TFlops單精度運(yùn)算性能,利用GPUBoost技術(shù)可達(dá)到單精度21.2TFlops運(yùn)算性能。

(2)NVLink技術(shù)。NVLink能夠支持多GPU間或CPU與GPU間的通信,能夠提供相當(dāng)于PCIe帶寬5倍的通信速率,雙向通信帶寬可達(dá)160GB/s。

當(dāng)NVLink被用來(lái)進(jìn)行CPU和GPU通信時(shí),前提條件就是CPU和GPU都支持NVLink技術(shù)。目前只有NVIDIA公司的部分GPU和IBM公司的Power8CPU支持NVLink技術(shù)。當(dāng)CPU和一個(gè)GPU利用NVLink技術(shù)連接時(shí),如圖9.31(a)所示,有4條NVLink連接,可以提供160GB/s的雙向通信帶寬。而當(dāng)CPU連接兩個(gè)GPU時(shí),如圖9.31(b)所示,每個(gè)GPU僅能提供兩條NVLink與CPU相連,GPU間有兩個(gè)NVLink連接,故雙向通信帶寬僅為80GB/s。圖9.31NVLink連接CPU和GPU

(3)HBM2堆疊內(nèi)存。P100首次在GPU中引入HBM2高速GPU存儲(chǔ)架構(gòu),訪存帶寬同比增長(zhǎng)了3倍,最高可以達(dá)到720GB/s。在Pascal架構(gòu)中,一個(gè)存儲(chǔ)控制器的訪存位寬為4096位,而Kepler架構(gòu)和Maxwell架構(gòu)中的GDDR5位寬為384位。

(4)統(tǒng)一存儲(chǔ)空間。P100首次在GPU中引入統(tǒng)一存儲(chǔ)空間,統(tǒng)一存儲(chǔ)空間提供了CPU和GPU存儲(chǔ)的統(tǒng)一地址空間(49位尋址空間,可容納512TB虛擬存儲(chǔ)空間)。利用統(tǒng)一存儲(chǔ),程序員無(wú)須關(guān)心數(shù)據(jù)存儲(chǔ)位置和設(shè)備間的通信。

(5)計(jì)算搶占。Pascal架構(gòu)允許計(jì)算任務(wù)在指令集被中斷,將程序上下文交換到DRAM,其他程序被調(diào)入和執(zhí)行。計(jì)算搶占技術(shù)能夠避免長(zhǎng)時(shí)間運(yùn)行的應(yīng)用獨(dú)占系統(tǒng)或發(fā)生運(yùn)行超時(shí),這樣,在與類似于交互式圖形或交互式調(diào)試任務(wù)進(jìn)行協(xié)同運(yùn)行時(shí),允許這類長(zhǎng)時(shí)運(yùn)行的應(yīng)用可以長(zhǎng)時(shí)間運(yùn)行來(lái)處理大規(guī)模數(shù)據(jù)或等待其他條件。另外,計(jì)算搶占可以打破kernel函數(shù)運(yùn)行時(shí)間限制,GPU程序在開(kāi)發(fā)時(shí)將不必考慮kernel函數(shù)運(yùn)行是否會(huì)超時(shí)。

(6)原子操作擴(kuò)展。Pascal架構(gòu)對(duì)原子操作進(jìn)行了擴(kuò)展,增加了64位數(shù)據(jù)的支持,包括64位的長(zhǎng)整型數(shù)據(jù)和雙精度浮點(diǎn)型數(shù)據(jù)。

9.6.3GPU編程方式

與傳統(tǒng)CPU將芯片大部分區(qū)域用于片上緩存不同,GPU將芯片絕大部分區(qū)域用于芯片邏輯,這導(dǎo)致現(xiàn)代GPU會(huì)有成百上千個(gè)內(nèi)核。為充分利用這些計(jì)算資源,必須為每個(gè)內(nèi)核創(chuàng)建一個(gè)線程。為了隱藏訪存延遲,甚至?xí)槊總€(gè)內(nèi)核創(chuàng)建多個(gè)線程。這就要求必須轉(zhuǎn)換原有的編程習(xí)慣。因?yàn)閺膸讉€(gè)線程到幾千個(gè)線程需要不同的方式來(lái)分配和處理任務(wù)負(fù)載。

1.圖形學(xué)API編程

最早的GPU只能執(zhí)行固定的幾類操作,沒(méi)有可編程的概念。利用DirectX和OpenGL等圖形API進(jìn)行程序映射時(shí),需要將計(jì)算的科學(xué)問(wèn)題轉(zhuǎn)換為圖形處理問(wèn)題,然后調(diào)用相應(yīng)的圖形處理接口完成計(jì)算,即所謂的可編程著色器(ProgrammableShader)。其后出現(xiàn)了相對(duì)高級(jí)的著色器語(yǔ)言(ShaderLanguage),例如,基于DirectX的HLSL和基于OpenGL后端的GLSL,以及同時(shí)支持DirectX和OpenGL的Cg。

2.Brook源到源編譯器

該編譯器由斯坦福大學(xué)的LanBuck等人在2003年開(kāi)發(fā)。Brook是對(duì)ANSIC的擴(kuò)展,是一個(gè)基于Cg的源到源編譯器,可以將類C語(yǔ)言的BrookC語(yǔ)言通過(guò)Brcc編譯器編譯為Cg代碼,很好地掩藏了圖形學(xué)API的實(shí)現(xiàn)細(xì)節(jié),大大簡(jiǎn)化GPU程序開(kāi)發(fā)過(guò)程。由于早期Brook只能使用像素著色器運(yùn)算,且缺乏有效的數(shù)據(jù)通信機(jī)制,導(dǎo)致效率低下。

3.Brook+

AMD/ATI在Brook基礎(chǔ)上,結(jié)合GPU計(jì)算抽象層(ComputeAbstractionLayer,CAL)推出Brook+。Brook+利用流與內(nèi)核的概念,在編程指定流和內(nèi)核后,由編譯器完成流數(shù)據(jù)和GPU的通信,運(yùn)行時(shí)自動(dòng)加載內(nèi)核到GPU執(zhí)行。內(nèi)核程序再編譯為AMD流處理器設(shè)備代碼IL,運(yùn)行時(shí)由CAL執(zhí)行。Brook+相比Brook有了巨大改進(jìn),但仍存在數(shù)據(jù)傳輸和流程序優(yōu)化困難的缺陷。AMD/ATI公司的StreamSDK中采用了Brook+作為高級(jí)開(kāi)發(fā)語(yǔ)言,用于AMD的Firestream系列GPU編程開(kāi)發(fā)。但目前StreamSDK和Brook+都已棄用,AMD產(chǎn)品主要以支持OpenCL為主。

4.CUDA

CUDA由NVIDIA在2007年發(fā)布,無(wú)須圖形學(xué)API,采用類C語(yǔ)言,開(kāi)發(fā)簡(jiǎn)單,到目前已發(fā)布CUDA10.0。CUDA支持C/C++、FORTRAN語(yǔ)言的擴(kuò)展,提供了豐富的高性能數(shù)學(xué)函數(shù)庫(kù),比如CUBLAS、CUFFT、CUDNN等。CUDA定義了GPU上執(zhí)行的數(shù)據(jù)和核函數(shù),通過(guò)運(yùn)行時(shí)API或設(shè)備API進(jìn)行設(shè)備和數(shù)據(jù)管理。CUDA結(jié)合GPU底層體系結(jié)構(gòu)特性,為用戶提供更底層控制,程序優(yōu)化具有巨大優(yōu)勢(shì)。

5.OpenCL

OpenCL是第一個(gè)面向異構(gòu)系統(tǒng)的通用并行編程標(biāo)準(zhǔn),也是一個(gè)統(tǒng)一的編程環(huán)境。OpenCL最初由蘋(píng)果公司提出,后由KhronosGroup發(fā)布并制定OpenCL行業(yè)規(guī)范,NVIDIA、Intel、AMD等IT巨頭均已支持OpenCL。OpenCL并行架構(gòu)包含宿主機(jī)和若干OpenCL設(shè)備,宿主機(jī)與OpenCL設(shè)備互連并整合為一個(gè)統(tǒng)一的并行平臺(tái),同時(shí)為程序提供API和運(yùn)行庫(kù)。主流的OFPGA和IntelXeonPhi等。penCL設(shè)備包括多核CPU、GPU、DSP(數(shù)字信號(hào)處理器)、

6.OpenACC

OpenACC最早由PGI公司提出并實(shí)現(xiàn),后被NVIDIA公司收購(gòu)。類似于OpenMP,OpenACC提供了一系列編譯指導(dǎo)指令,通過(guò)在程序并行區(qū)域外指定編譯指導(dǎo)語(yǔ)句,然后由編譯器對(duì)并行區(qū)域內(nèi)代碼進(jìn)行分析,編譯為目標(biāo)平臺(tái)上的源代碼,這是一種源到源的轉(zhuǎn)換。OpenACC隱藏了異構(gòu)系統(tǒng)主機(jī)端和設(shè)備端之間數(shù)據(jù)傳輸和執(zhí)行調(diào)度等細(xì)節(jié),大大簡(jiǎn)化了異構(gòu)編程。OpenACC的執(zhí)行模型包括gang、worker和vector三級(jí)并行結(jié)構(gòu)。

7.OpenMP4.5

OpenMP4.5是對(duì)標(biāo)準(zhǔn)共享存儲(chǔ)編程模型OpenMP的擴(kuò)展,擴(kuò)展內(nèi)容主要是支持異構(gòu)計(jì)算。通過(guò)在指導(dǎo)命令層指定數(shù)據(jù)傳輸和描述加速任務(wù)區(qū)來(lái)實(shí)現(xiàn)異構(gòu)計(jì)算。目前Intel的ICC編譯器已經(jīng)有了OpenMP4.0的實(shí)現(xiàn),其加速器主要支持IntelXeonPhi。NVIDIA預(yù)計(jì)在OpenMP4.5支持GPU加速,GCC已有初步的支持。另外,若有多個(gè)GPU或在GPU集群中,可以聯(lián)用OpenMP或MPI與GPU編程方法(比如CUDA)進(jìn)行混合編程。MPI+OpenMP+CUDA是GPU集群上主流的混合編程模型。

CUDA編程模型為全局串行局部并行(GloballySequentialLocallyParallel,GSLP)模型,如圖9.32所示。在CUDA編程模型中引入了主機(jī)端和設(shè)備端的概念,其中CPU作為主機(jī)端,GPU作為設(shè)備端,主機(jī)端僅有一個(gè),而設(shè)備端可以有多個(gè)。CPU負(fù)責(zé)邏輯處理和運(yùn)算量少的計(jì)算,GPU負(fù)責(zé)運(yùn)算量大的并行計(jì)算。完整的CUDA程序包括主機(jī)端和設(shè)備端兩部分代碼,主機(jī)端代碼在CPU上執(zhí)行,設(shè)備端代碼(又稱kernel函數(shù))運(yùn)行在GPU上。為了充分利用GPU,應(yīng)用程序必須分解為能夠并發(fā)執(zhí)行的大量線程,GPU調(diào)度器能夠以最小的交換開(kāi)銷并根據(jù)實(shí)際設(shè)備的執(zhí)行能力高效地執(zhí)行這些線程。圖9.32CUDA全局串行局部并行編程模型

CUDA運(yùn)行時(shí)提供了豐富的函數(shù),功能涉及設(shè)備管理、存儲(chǔ)管理、數(shù)據(jù)傳輸、線程管理、流管理、事件管理、紋理管理、執(zhí)行控制、與OpenGL和Direct3D互操作等。一個(gè)完整的CUDA代碼需要七個(gè)關(guān)鍵步驟,特別是多GPU情況下;若是單GPU,可省略為五個(gè)步驟。下面簡(jiǎn)單描述CUDA代碼的七個(gè)關(guān)鍵步驟:

學(xué)習(xí)使用CUDA庫(kù)函數(shù)能夠幫助GPU程序員快速開(kāi)發(fā)高性能的CUDA程序。針對(duì)向量加法運(yùn)算,在CUBLAS庫(kù)函數(shù)中的cublas<t>axpy()函數(shù)可以實(shí)現(xiàn)向量加法的功能。函數(shù)調(diào)用時(shí)要注意以下幾點(diǎn):

(1)包含頭文件cublas_v2.h。

(2)定義cublasHandle_t句柄。

(3)主機(jī)端到設(shè)備端數(shù)據(jù)傳輸函數(shù)是cublasSetVector(),設(shè)備端到主機(jī)端數(shù)據(jù)傳輸函數(shù)是cublasGetVector()。

(4)cublasSaxcpy()函數(shù)調(diào)用時(shí)為cublasSaxcpy_v2()。

9.6.4多線程多處理器體系結(jié)構(gòu)

GPU是一種多線程多處理器體系結(jié)構(gòu)的協(xié)處理器。GPU的每一個(gè)處理器可以有效地執(zhí)行多線程的程序。高性能計(jì)算采用GPU一般會(huì)有幾個(gè)到幾十個(gè)不等的多線程處理器。GPU采用多線程處理器而不