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

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

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

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

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

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

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

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

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

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

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

1.片內(nèi)并行

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

2.協(xié)處理器

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

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

3.多處理器

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

4.多計算機

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

5.云計算

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

(4)拓撲結(jié)構(gòu)也極大地影響著網(wǎng)絡(luò)的實現(xiàn)成本和整個計算機系統(tǒng)的性能價格比。

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

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

1.圖

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

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

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

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

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

2.互連函數(shù)

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

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

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

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

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

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

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

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

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

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

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

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

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

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

解因為CPU2與內(nèi)存模塊0之間的連線交叉點為閉合的交叉開關(guān)(黑圈),可表示為

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

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

①它是無阻塞網(wǎng)絡(luò),即節(jié)點之間不會因為某些交叉點或者鏈路被占用而無法建立連接(僅在有兩個以上節(jié)點試圖向同一個目的節(jié)點發(fā)信息時才可能在目的節(jié)點接收鏈路上發(fā)生阻塞);

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

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

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

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

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

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

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

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

(2)交換開關(guān)之間采用何種拓撲連接。

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

最常選用的交換開關(guān)是2×2交換開關(guān),因為它結(jié)構(gòu)簡單,容易控制,成本低。級間連接常用的拓撲有全混洗、蝶形、縱橫交叉、立方體連接等。對交換開關(guān)的控制一般有以下四種方式:

(1)整體控制:所有交換開關(guān)使用同一個控制信號控制。

(2)級控制:同一級交換開關(guān)使用同一個控制信號控制。

(3)單元級控制:每個交換開關(guān)分別控制。

(4)部分級控制:同一級中部分交換開關(guān)使用同一個控制信號控制。

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

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

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

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

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

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

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

圖9.10兩種Bene?拓撲結(jié)構(gòu)

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

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

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

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

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

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

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

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

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

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

用于分布式交換網(wǎng)絡(luò)的拓撲結(jié)構(gòu)與集中式交換網(wǎng)絡(luò)有幾點不同:

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

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

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

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

1.線性網(wǎng)絡(luò)拓撲結(jié)構(gòu)

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

2.環(huán)形網(wǎng)絡(luò)拓撲結(jié)構(gòu)

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

3.網(wǎng)格形網(wǎng)絡(luò)拓撲結(jié)構(gòu)

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

4.超立方體網(wǎng)絡(luò)拓撲結(jié)構(gòu)

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

9.4向量處理機

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

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

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

(2)向量指令使結(jié)果的向量元素計算互不相關(guān),因此硬件在一條向量指令執(zhí)行期內(nèi)不必檢查數(shù)據(jù)相關(guān)(冒險)?？梢岳貌⑿械墓δ懿考嚵?、單一的深度流水的功能部件或者并行和流水的功能部件的組合來計算向量的各元素。

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

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

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

所以,向量運算比具有相同數(shù)據(jù)項的標量運算更快。向量處理機主要用于大型科學和工程計算中。

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

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

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

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

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

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

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

(4)標量寄存器組。標量寄存器也能夠為向量功能部件提供輸入數(shù)據(jù),同時計算傳遞給向量Load-Store部件的地址。它們是常規(guī)32個通用寄存器和32個MIPS浮點寄存器。標量值從標量寄存器文件讀出,然后鎖存在向量功能部件的一個輸入端。圖9.19計算DAXPY算法

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

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

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

圖9.20CrayMSP模塊

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

1)寫直達協(xié)議

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

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

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

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

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

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

(2)僅當需要替換該Cache塊時,將該Cache塊寫回共享存儲器。

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

(2)使用Load和Store指令訪問遠程共享存儲器。

(3)訪問遠程共享存儲器比訪問本地共享存儲器速度慢。

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

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

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

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

9.6.1簡介隨著圖形顯示器的大規(guī)模應(yīng)用以及計算機游戲行業(yè)(包括PC和專用游戲機)的快速增長,傳統(tǒng)的微處理器對于圖形處理力不從心,摩爾定律又限制了微處理器可用的晶體管數(shù)量,因此許多公司加大對圖形處理器的投資,也使得圖形處理器的發(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),通過將訪存控制器集成到CPU芯片上減少訪存延遲。圖9.27(c)為AMD加速處理單元(AcceleratedProcessorUint,APU),它將CPU和GPU集成到同一芯片上,這是一個巨大的突破,特別是隨著異構(gòu)統(tǒng)一內(nèi)存訪問技術(shù)的引入,不但統(tǒng)一了CPU和GPU的內(nèi)存空間,而且還在兩者間保證了緩存的一致性。圖9.27CPU-GPU架構(gòu)

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

(1)浮點計算能力強。

(2)超高性價比。

(3)綠色功耗比。

(4)普及度廣。

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

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

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

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

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

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

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

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

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

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

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

(10)存取單元(Load/StoreUnits,LD/ST):負責全局內(nèi)存、共享內(nèi)存、常量內(nèi)存、紋理加載等存儲器訪問操作。

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

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

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

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

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

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

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

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

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

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

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

9.6.3GPU編程方式

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

1.圖形學API編程

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

2.Brook源到源編譯器

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

3.Brook+

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

4.CUDA

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

5.OpenCL

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

6.OpenACC

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

7.OpenMP4.5

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

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

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

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

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

(2)定義cublasHandle_t句柄。

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

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

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

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