最近一段時(shí)間，有關(guān)超算的話題成為熱門，一時(shí)間大家都開始討論超算，然而，筆者發(fā)現(xiàn)在所有這些討論中，從沒有在任何時(shí)間任何地點(diǎn)發(fā)現(xiàn)任何人問出就連小學(xué)生都經(jīng)常問的問題：超算到底是怎么算的？不得不說是一件可悲的事情。

　　【并行計(jì)算怎么算】

　　我們知道，單個(gè)CPU核心只能串行計(jì)算，也就是一條一條的把機(jī)器指令讀出來執(zhí)行。要理解的一點(diǎn)是，串行執(zhí)行并不表示指令順序執(zhí)行，跳轉(zhuǎn)指令可以讓CPU跳到其他地址執(zhí)行，但是整個(gè)過程CPU只在執(zhí)行一個(gè)單一的指令流，也就是一個(gè)線程，thread。某個(gè)線程完成某種任務(wù)，而線程對(duì)應(yīng)的代碼中的多個(gè)函數(shù)又各自分工完成對(duì)應(yīng)的工序。要想同時(shí)執(zhí)行多個(gè)線程的話，有兩個(gè)辦法，一個(gè)是增設(shè)額外的一個(gè)或者多個(gè)CPU，這樣，在時(shí)間上可以做到Parallel/并行，同一時(shí)刻有多個(gè)任務(wù)同時(shí)執(zhí)行；另一種辦法則是讓單個(gè)CPU執(zhí)行一段時(shí)間線程1，然后再強(qiáng)行跳轉(zhuǎn)到線程2執(zhí)行一段時(shí)間，然后再跳回到線程1，這樣就可以實(shí)現(xiàn)多個(gè)線程的Concurrency/并發(fā)，但是卻不是并行，因?yàn)橥粋€(gè)時(shí)刻還是只有一個(gè)線程在執(zhí)行，只不過每個(gè)線程執(zhí)行的時(shí)間非常短，一般比如10ms的時(shí)間，就會(huì)跳轉(zhuǎn)到其他線程執(zhí)行，這樣從表面看來，一段時(shí)間內(nèi)，多個(gè)線程似乎是“同時(shí)”執(zhí)行的。

　　前者的方式看上去性能更高，但是其有2個(gè)慘痛的代價(jià)，第一個(gè)是線程之間的同步，第二個(gè)是緩存一致性。如果多個(gè)線程運(yùn)行在同一個(gè)核心上，那么它們只能一個(gè)接一個(gè)的執(zhí)行，執(zhí)行線程1時(shí)線程2不可能得到執(zhí)行，如果線程1和線程2要操作同一個(gè)變量，那么就輪流操作，不會(huì)有問題。但是多個(gè)線程運(yùn)行在不同的核心上，事情就發(fā)生很大變化了，比如有兩個(gè)線程，都需要操作某個(gè)變量，比如同時(shí)運(yùn)行a=a+1這個(gè)邏輯，期望結(jié)果是線程1對(duì)a加了1，線程2要在線程1輸出結(jié)果的基礎(chǔ)上繼續(xù)+1，而由于這兩個(gè)線程運(yùn)行在兩個(gè)獨(dú)立核心上彼此之間沒有協(xié)調(diào)，可能導(dǎo)致線程1讀到的a的初始值0，加1之后還沒來得及將最新結(jié)果更改到a所在的地址之前，線程2也讀到了a的初始值0，加1之后也嘗試寫入同樣的地址，最后a的結(jié)果是1，而不是期望中的2。解決辦法則是對(duì)變量a加互斥鎖，當(dāng)某個(gè)線程操作a之前，先將鎖（也是個(gè)變量）置為1，其他線程不斷的掃描鎖是不是已經(jīng)被置為0，如果是1則表示其他人正在操作a，如果是0則表示其他人已經(jīng)釋放了，那么其將鎖改為1，也就是鎖上，自己操作a，此時(shí)讀到的a就會(huì)是被其他線程更新之后的最新數(shù)據(jù)了。這個(gè)過程叫做Consistency。所以，如果多個(gè)線程之間完全獨(dú)立各干各的，沒有任何交互，這是最理想的場景，這就像多臺(tái)無須聯(lián)網(wǎng)的獨(dú)立的計(jì)算機(jī)各干各的一樣，只不過共用了CPU和內(nèi)存。

　　然而，如果使用了緩存，又使用了共享變量，事情又變得復(fù)雜了。線程1所在的核心1搶到變量a的鎖之后會(huì)將a的內(nèi)容緩存到核心1的緩存中，更新了a內(nèi)容之后，該更新也依然留在緩存中而不是被flush到主存。其釋放鎖之后，線程2搶到a的鎖并將a讀入核心2的緩存，此時(shí)如果不做任何處理，核心2從主存中讀到的將是a的舊內(nèi)容，從而計(jì)算出錯(cuò)?？梢钥吹剑幢闶鞘褂昧随i來保證Consistency，也無法避免緩存所帶來的一致性問題，后者則被稱為Coherency。

　　Consistency由軟件來負(fù)責(zé)，而Coherency則要由硬件來負(fù)責(zé)保證，具體做法是將每一筆數(shù)據(jù)更新同步廣播出去給所有其他核心/CPU，將它們緩存中的舊內(nèi)容作廢，收到其他所有核心的回應(yīng)之后，該更新才被認(rèn)為成功。所以核心/CPU之間需要一個(gè)超低時(shí)延的網(wǎng)絡(luò)用于承載這個(gè)廣播。這個(gè)過程對(duì)軟件完全透明。除了需要廣播作廢外，當(dāng)其他核心需要訪問該變量時(shí)，擁有該變量最新內(nèi)容的核心必須做出應(yīng)答將該內(nèi)容推送到發(fā)出訪問請(qǐng)求的核心。

　　在很早期的SMP/UMA架構(gòu)下，由于那時(shí)的SMP總線本身就是一個(gè)廣播域，任何核心的訪存請(qǐng)求都會(huì)被其他所有核心收聽到，包括更新了某個(gè)地址、讀取某個(gè)地址，這樣很天然的可以實(shí)現(xiàn)Coherency，比如，當(dāng)某個(gè)核心更新了某個(gè)地址之后，其他核心后臺(tái)默默的收聽（或者說嗅探，Snoop），并在自己緩存中查詢自己有沒有緩存這個(gè)地址的內(nèi)容，有則作廢無則不動(dòng)作。當(dāng)某個(gè)核心發(fā)起對(duì)某個(gè)地址讀的時(shí)候，其他核心收聽之后也默默的搜索自己的緩存看看是否有該內(nèi)容最新版本，有則使用特殊的信號(hào)搶占總線并壓制主存控制器對(duì)總線的搶占，將數(shù)據(jù)返回到總線上，與此同時(shí)主存控制器也收聽該內(nèi)容并將該內(nèi)容同步更新到該地址在主存中的副本。此時(shí)，該地址將擁有三個(gè)副本，分別位于：之前緩存它的那個(gè)核心的緩存、剛剛讀它的那個(gè)核心的緩存、主存，而且內(nèi)容一致。如果此時(shí)之前緩存它的那個(gè)核心再次發(fā)起讀操作，就沒有必要將讀請(qǐng)求發(fā)送到總線上，而浪費(fèi)電，同時(shí)也浪費(fèi)其他核心的搜索運(yùn)算耗費(fèi)的電能以及對(duì)其他正常緩存訪問的搶占。所以人們想了個(gè)辦法，每個(gè)緩存條目（Cache Line，緩存行）增加一個(gè)字段，專門用來描述”該緩存當(dāng)前處于什么狀態(tài)“，上述狀態(tài)稱為Share態(tài)，而如果有人更新了某個(gè)地址，其他核心嗅探到之后，便將自己緩存里這份內(nèi)容改為”Invalid“態(tài)，而剛剛更新內(nèi)容的那個(gè)核心里緩存的該條目被改為”Modified“態(tài)，Invalid態(tài)的條目已經(jīng)作廢，再讀就得走總線，M態(tài)的條目可以直接讀，因?yàn)榇藭r(shí)沒有其它人有比你新的內(nèi)容了。如果加電之后某個(gè)核心第一個(gè)搶到總線并發(fā)起該地址的讀，則讀入之后該條目就是Exclusive態(tài)，因?yàn)橹挥兴粋€(gè)人緩存了該條目，當(dāng)另外核心再發(fā)起讀之后，該核心嗅探到這個(gè)事件，于是將自己緩存里的該條目發(fā)送給剛才發(fā)起讀的核心，那個(gè)核心從而知道其他核心也有該內(nèi)容，于是兩個(gè)核心一起將自己本地的該條目改為Share態(tài)，這兩個(gè)核心中任何一個(gè)如果再發(fā)起該地址的讀，就不用走總線了，直接緩存命中?？梢钥吹剑?dāng)某個(gè)核心需要訪問的數(shù)據(jù)在其他核心的緩存中時(shí)，硬件會(huì)自動(dòng)傳遞這份數(shù)據(jù)，軟件根本無需關(guān)心。所以多個(gè)線程之間引用共享變量的時(shí)候，直接引用即可。

　　上述方式被稱為MESI協(xié)議，其目的是為了提升效率，不需要每一筆訪問都走外部總線。后來過渡到NUMA架構(gòu)之后，NUMA是通過一個(gè)分布式交換網(wǎng)絡(luò)來廣播同步這些消息以及進(jìn)行變量內(nèi)容傳送的，由于該網(wǎng)絡(luò)并非一跳直達(dá)的廣播網(wǎng)絡(luò)，所以過濾不必要的廣播就更加重要了。不同的CPU廠商有不同的方式，MESI協(xié)議也有不少變種，比如MESIF等等。另外，由于失去了天然的總線嗅探機(jī)制，如果某個(gè)緩存行處于Invalid態(tài)，讀取該緩存行之前硬件需要發(fā)出Probe操作，探尋，主動(dòng)廣播這個(gè)Probe請(qǐng)求給所有核心的緩存控制器，緩存了該行的緩存控制器會(huì)返回最新數(shù)據(jù)并將自己該行的狀態(tài)改為S態(tài)。如果要更新某行，該行本地處于E或者M(jìn)態(tài)則直接更新，處于S態(tài)則需要發(fā)出Probe請(qǐng)求作廢其他緩存中的該行?？傊?，當(dāng)MESI遇到NUMA，就是個(gè)非常復(fù)雜的狀態(tài)機(jī)，筆者就不繼續(xù)展開了，要想與筆者深聊NUMA和Cache Coherency，可接受預(yù)約面聊，前提是你得拿出筆者看得上的干貨來咱們互換一下。

　　【核心數(shù)能否再多點(diǎn)】

　　可以看到，增加核心數(shù)量并不是那么容易的事情，除非不使用緩存，所有核心吧所有的更新都寫到主存里。但是這樣做性能將會(huì)不可接受。有人問，把緩存也集中共享不就沒這么多事了么？的確，但是如果把緩存單獨(dú)放到某個(gè)地方，多個(gè)CPU芯片通過某種總線集中訪問該緩存，那么其總線速率一定不夠高，因?yàn)槠渥叩搅诵酒饷?，?dǎo)線長度變高，信號(hào)質(zhì)量就會(huì)變差。這個(gè)思路就不現(xiàn)實(shí)了。

　　所以，又得使用分布式緩存，又得要求核心數(shù)量越來越多的話，就得保證所有核心之間的網(wǎng)絡(luò)足夠高速才行，而核心數(shù)量越來越多，網(wǎng)絡(luò)的直徑就會(huì)越來越大，即便是MESI過濾，去效果也是有限的，廣播的時(shí)延隨著網(wǎng)絡(luò)規(guī)模的增大變得越來越高。所以，核心數(shù)量達(dá)到一定程度之后，緩存一致性問題就變成了整個(gè)系統(tǒng)擴(kuò)展性的瓶頸點(diǎn)所在。怎么辦？沒辦。反正，硬件是不會(huì)再給你保證一致性了，因?yàn)槿绻砂偕锨€(gè)核心的網(wǎng)絡(luò)，用硬件保證一致性得不償失。

　　此時(shí)，必須拋棄由硬件保證的緩存一致性，改為軟件自行解決。NoC（Network on Chip）方案就是在一個(gè)芯片中將幾十個(gè)上百個(gè)核心通過高速網(wǎng)絡(luò)連接起來但是卻不提供硬件緩存一致性，其網(wǎng)絡(luò)直徑很大。這種架構(gòu)天然適合各干各的，如果不是各干各的，必須傳遞更新后的變量的話，需要由軟件自行向位于每個(gè)核心前端的NoC網(wǎng)絡(luò)控制器發(fā)送消息+目標(biāo)節(jié)點(diǎn)地址并傳遞到對(duì)方，對(duì)方通過底層驅(qū)動(dòng)+協(xié)議棧接收該變量并傳遞給其本地程序，這已經(jīng)是赤裸裸的程序控制網(wǎng)絡(luò)通信了，其實(shí)NUMA已經(jīng)是這樣了，只不過其網(wǎng)絡(luò)通信程序跑在硬件微碼或者硬狀態(tài)機(jī)中，且該狀態(tài)機(jī)可直接接收訪存請(qǐng)求并將其通過網(wǎng)絡(luò)發(fā)送從而對(duì)軟件透明。

　　另外，NoC架構(gòu)的CPU很少會(huì)被設(shè)計(jì)為共享內(nèi)存架構(gòu)，因?yàn)榇藭r(shí)主存也是通過主存控制器接入NoC，由于NoC時(shí)延過大，每一筆訪存請(qǐng)求又是同步的，將代碼直接放到主存，性能將會(huì)非常差，所以NoC上的RAM主要用于所有核心之間的最后一層共享緩存了，只不過是可尋址的緩存，由軟件而不是硬件來管理。NoC架構(gòu)下每個(gè)核心內(nèi)部一般會(huì)有幾百KB的SRAM可尋址空間，代碼則運(yùn)行在這里。外部主存可以使用虛擬驅(qū)動(dòng)映射成某個(gè)帶隊(duì)列的設(shè)備，異步讀寫。也有支持直接映射到核心地址空間的但是訪問性能會(huì)很差，所以鮮有?！?/p>

　　【再多點(diǎn)！】　

　　我說，能否讓核心數(shù)量再多點(diǎn)？筆者大學(xué)是應(yīng)用化學(xué)專業(yè)，但是筆者學(xué)到最后已經(jīng)完全對(duì)化學(xué)失去了興趣，尤其是看到薛定諤老哥們之后。記得某室友的畢業(yè)設(shè)計(jì)就是用程序來算分子結(jié)構(gòu)，一臺(tái)破pc，奔四的，導(dǎo)師/師兄給了個(gè)程序說：按照步驟輸入?yún)?shù)，點(diǎn)確定，貼個(gè)條“別關(guān)機(jī)，計(jì)算中！“，行了。一個(gè)月以后，來看結(jié)果。的確，不夠用啊，如果一分鐘能算完，這哥們就可以更安逸的玩暗黑2和看玄幻了。錢啊，學(xué)費(fèi)啊，就特么干這個(gè)。當(dāng)然，筆者做的更讓人寒心，當(dāng)時(shí)筆者是攪和化合物，分析天平稱點(diǎn)粉末，弄點(diǎn)液體放燒瓶里加熱攪拌，一晚上第二天結(jié)晶，拿去做x光衍射分析，寫論文，當(dāng)時(shí)都不知道自己在干什么，為什么干。這比那哥們更耗錢，人家只耗電，筆者這一勺子下去，就是幾千塊！撒地上一點(diǎn)就是幾百塊?。“?，內(nèi)疚?。?/p>

　　跑題了。核心數(shù)量再多，一個(gè)芯片真就搞不定了，NoC也白搭。必須用多個(gè)芯片，也不行，幾萬個(gè)CPU芯片怎么給它弄到一起？那就不能NoC了，得眼睛看得見的網(wǎng)絡(luò)了，包括網(wǎng)卡、網(wǎng)線、交換機(jī)。這么大范圍的網(wǎng)絡(luò)，其速率相比NoC又降低一個(gè)檔次，比如萬兆以太、Infiniband等這種級(jí)別了。也就是說，這種規(guī)模之下，必須使用多臺(tái)獨(dú)立的機(jī)器來搭建，也就是所謂的超算集群。

　　哦？一堆單獨(dú)的機(jī)器用網(wǎng)絡(luò)連起來就是超級(jí)計(jì)算機(jī)？難道不應(yīng)該是科幻片里那些超腦類腦神秘機(jī)器么？哈哈，沒那么科幻了，兄弟?？梢赃@么說，整個(gè)Internet上的計(jì)算機(jī)也天然組成了一個(gè)超算集群，只要Internet用戶同意，就可以用來計(jì)算。比如SETI@Home項(xiàng)目就是利用所有Internet上的計(jì)算機(jī)各自下載一部分?jǐn)?shù)據(jù)然后用同樣的方式去分析然后返回結(jié)果，其將程序作為一個(gè)屏保程序，在離開或者空閑的時(shí)候，便后臺(tái)啟動(dòng)計(jì)算，屏保結(jié)束則自動(dòng)停止。當(dāng)然，這個(gè)計(jì)算過程中基本不會(huì)有網(wǎng)絡(luò)通信，多個(gè)部分之間各算個(gè)的。

　　使用了外部網(wǎng)絡(luò)的話，多個(gè)不同機(jī)器上的線程之間就不能夠共享內(nèi)存來通信或者共同引用或處理某個(gè)變量或數(shù)據(jù)結(jié)構(gòu)。如果需要通信或者共享變量，就得徹底使用這些外部網(wǎng)絡(luò)來傳遞了。比如，利用TCP/IP，RDMA over IB等等。這么復(fù)雜了？老子之前寫程序都是直接聲明變量和數(shù)據(jù)結(jié)構(gòu)，直接引用的，充其量加個(gè)鎖，你現(xiàn)在讓老子每次引用的時(shí)候要調(diào)用TCP/IP發(fā)個(gè)包給對(duì)方，再接收包，才能拿到？老子不干。。

　　不干不行啊，必須干，你不干就別用幾萬個(gè)核，回去舒舒服服用你的8路服務(wù)器，或者小型機(jī)去，那個(gè)舒服，但是就是一算一個(gè)月，爽不？酸爽。那好吧，老子接受，但是底層這么多網(wǎng)絡(luò)，老子不懂TCP/IP怎么調(diào)用，更不懂什么RDMA，幾千上萬臺(tái)機(jī)器，老子光記錄IP地址就得記多少個(gè)。。你說咋辦吧。

　　【MPI是干啥的】

　　鑒于上面這位老兄的顧慮，人們開發(fā)了一種叫做Message Passing Interface，簡稱MPI的函數(shù)庫。來來來，這位客官，您要把什么東西發(fā)送給誰，敬請(qǐng)吩咐~~。嗯~~~？好，我這有個(gè)數(shù)組a［100］，現(xiàn)在想把a(bǔ)［0］到a［9］分別發(fā)給運(yùn)行在1~9號(hào)節(jié)點(diǎn)上的各自進(jìn)程。片刻后，“上~~菜！”。我去，這么快就完成了？

　　上述場景成為MPI_Scatter場景，某個(gè)進(jìn)程將一組數(shù)據(jù)一個(gè)一個(gè)的散播給多個(gè)進(jìn)程。對(duì)應(yīng)的函數(shù)為MPI_Scatter （&sendbuf，sendcnt，sendtype，&recvbuf，recvcnt，recvtype，root，comm），每個(gè)進(jìn)程的代碼中該處加上一句，就可以了，該收的收，該發(fā)的發(fā)，該函數(shù)為同步阻塞調(diào)用，執(zhí)行到這一步時(shí)發(fā)送方發(fā)出數(shù)據(jù)，接收方等待接收，搞完之后大家再繼續(xù)往下走。

　　如果是共享內(nèi)存架構(gòu)，進(jìn)程在代碼里可以直接引用a［i］，根本不需要接收/發(fā)送。比如int b=a［9］+1。該函數(shù)底層也是調(diào)用本機(jī)的網(wǎng)絡(luò)協(xié)議棧、設(shè)備驅(qū)動(dòng)從而將消息封包傳遞給其他機(jī)器的。

　　同理，還有如下這些常用函數(shù)：MPI_Send， MPI_ISend， MPI_Recv， MPI_IRecv等用于點(diǎn)對(duì)點(diǎn)數(shù)據(jù)傳遞的；MPI_Bcase， MPI_Gather， MPI_Reduce， MPI_Barrier等用于集合通信的等，后者統(tǒng)稱為Collective類MPI通信。

　　利用MPI框架編寫好對(duì)應(yīng)的程序之后，需要使用mpirun腳本將程序load到集群里的每個(gè)節(jié)點(diǎn)上，具體如何load需要預(yù)先寫好一份配置文件，大家有興趣可以自行研究，筆者這里就不再多說了。

　　【浪潮Caffe MPI版本】

　　浪潮則是國內(nèi)HPC領(lǐng)域的領(lǐng)軍廠商，有一系列的軟硬件解決方案，比如高密度刀片、GPU一體機(jī)、深度學(xué)習(xí)一體機(jī)、SmartRack整機(jī)柜服務(wù)器、InCloudRack整機(jī)柜服務(wù)器等等。

　　在HPC生態(tài)環(huán)境和軟件方面，浪潮專門為Caffe深度學(xué)習(xí)平臺(tái)開發(fā)了對(duì)應(yīng)的MPI版本（開源）；提供高性能計(jì)算服務(wù)平臺(tái)ClusterEngine，可實(shí)現(xiàn)集群監(jiān)控，故障報(bào)警，系統(tǒng)管理，作業(yè)提交管理及調(diào)度，支持提供優(yōu)先級(jí)，公平共享，資源限制，資源回填等多種調(diào)度策略，支持HPC應(yīng)用集成，記賬管理，統(tǒng)計(jì)分析等功能；

　　以及集群診斷調(diào)優(yōu)工具Teye。天眼是一個(gè)高性能應(yīng)用特征監(jiān)控分析系統(tǒng)，面向大規(guī)模HPC集群系統(tǒng)，用于提取應(yīng)用程序?qū)嘿Y源的使用情況，并實(shí)時(shí)反映其運(yùn)行特征；可在現(xiàn)有硬件平臺(tái)基礎(chǔ)上，深度挖掘應(yīng)用程序的計(jì)算潛力；為診斷應(yīng)用程序瓶頸，改進(jìn)應(yīng)用算法，提高并行效率提供科學(xué)有力的指引，最終達(dá)到優(yōu)化系統(tǒng)資源利用率，提高系統(tǒng)計(jì)算性能的目標(biāo)。并提供熱點(diǎn)走勢圖、數(shù)據(jù)分布律、特征雷達(dá)圖等可視化監(jiān)控。

　　深度學(xué)習(xí)是近年來的熱點(diǎn)技術(shù)，筆者在這方面的知識(shí)為0。雖然不知道它是怎么做的，但是看上去真的很厲害。Caffe是伯克利發(fā)布的一款深度學(xué)習(xí)計(jì)算框架，廣泛用于圖像識(shí)別等深度學(xué)習(xí)領(lǐng)域。

　　浪潮發(fā)布的MPI集群版Caffe計(jì)算框架正是切中當(dāng)下深度學(xué)習(xí)的迫切需求，其基于伯克利的caffe框架進(jìn)行了MPI改造，讓原本只支持單機(jī)+GPU環(huán)境的Caffe框架完美支持了多機(jī)MPI集群環(huán)境，并且完全保留了caffe原有的特性。

　　浪潮的MPI版的Caffe計(jì)算框架已經(jīng)在某超級(jí)計(jì)算機(jī)上進(jìn)行部署并測試，結(jié)果顯示，在保證正確率相同的情況下，浪潮MPI Caffe在16個(gè)GPU上并行計(jì)算效率上提升13倍。

　　【GPU事半功倍】

　　超算又可以被稱為并行計(jì)算，用大量的計(jì)算核心堆砌。高端GPU里天然就有數(shù)千個(gè)計(jì)算核心，其原本專門用于加速計(jì)算圖形渲染過程，后來被人們用于通用計(jì)算領(lǐng)域。Nvidia GPU的典型架構(gòu)如下圖所示。

　　上圖只是多個(gè)Stream Multi Processor中的一個(gè)，整個(gè)GPU含有數(shù)以千計(jì)的Core。程序編寫人員需要通告GPU某個(gè)計(jì)算步驟要分成多少個(gè)線程以及如何劃分。GPU內(nèi)部會(huì)將32個(gè)線程為一組（也就是warp），輪流調(diào)度到SM中執(zhí)行。

　　比如如下對(duì)一個(gè)一維數(shù)組中的各項(xiàng)元素求平方然后再求和的過程：

　　對(duì)上述過程的計(jì)算加速很簡單，就是創(chuàng)建多個(gè)線程，裝載到多個(gè)核心上同時(shí)執(zhí)行，每個(gè)核心執(zhí)行這一百萬個(gè)元素中的一部分，形成部分和，最后再累加起來即可。如果形成一千個(gè)線程，每個(gè)線程各自求解其中一千個(gè)元素的平方和，然后再累加一千次即可，累加的過程也可以使用加法樹再次加速，比如再將這一千個(gè)部分和分成20個(gè)線程來加，每個(gè)線程加50個(gè)數(shù)，這樣同一時(shí)段內(nèi)總共累加70次即可完成，而不是累加一千次。上述過程可以采用OpenMP庫在單機(jī)共享內(nèi)存場景下來完成，OpenMP庫會(huì)自動(dòng)將這個(gè)for循環(huán)展開為多個(gè)線程計(jì)算。如果采用NVidia GPU來加速上述過程，則需要下列步驟：

　　Step1. 主程序在主機(jī)內(nèi)存中將數(shù)據(jù)準(zhǔn)備好，比如：int data［1048576］，表示一個(gè)含100萬個(gè)元素的一維數(shù)組，然后向其中填充對(duì)應(yīng)的待計(jì)算的數(shù)據(jù)。

　　Step2. 主程序調(diào)用NVidia提供的CUDA庫函數(shù)InitCUDA（）聲明當(dāng)前進(jìn)程需要與GPU通信調(diào)用其執(zhí)行計(jì)算任務(wù)，GPU側(cè)則做好相應(yīng)的上下文切換。

　　Step3. 主程序調(diào)用Nvidia所提供的CUDA庫函數(shù)cudaMalloc（ ），要求GPU在顯存中開辟對(duì)應(yīng)大小的空間用于存放第一步中生成的數(shù)據(jù)。

　　Step4. 主程序調(diào)用CUDA庫函數(shù)cudaMemcpy（ ）讓GPU從主存中對(duì)應(yīng)位置將數(shù)據(jù)拷貝到第二步開辟的顯存空間中。

　　Step5. 主程序采用 ”__global__ 函數(shù)名《《《塊數(shù)，線程數(shù)》》》 （參數(shù)） 的方式聲明讓CUDA庫自動(dòng)在GPU上生成多個(gè)線程來同時(shí)執(zhí)行該函數(shù)，該函數(shù)被CUDA成為“核函數(shù)”或者“Kernel”，意思就是該函數(shù)中才是整個(gè)計(jì)算的核心邏輯，會(huì)被GPU并行執(zhí)行，之前做的都是些準(zhǔn)備工作而已。在函數(shù)中也可以采用手動(dòng)方式明確指定由哪個(gè)線程處理哪一部分的的數(shù)據(jù)，具體采用將線程ID與變量相關(guān)聯(lián)的方式實(shí)現(xiàn)。《《《塊數(shù)，線程數(shù)》》》是CUDA中對(duì)線程的組織方式，多個(gè)線程組成塊，多個(gè)塊組成Grid，用三級(jí)描述來區(qū)分各個(gè)線程。調(diào)度則以32個(gè)線程為一組。

　　Step6. 主程序調(diào)用cudaMemcpy（ ）將第五步執(zhí)行完的結(jié)果從顯存拷貝到主存。

　　Step7. 主程序調(diào)用cudafree（ ）釋放之前分配的顯存。

　　Step8. 主程序進(jìn)行收尾工作，比如進(jìn)行部分和的累加，以及顯示出對(duì)應(yīng)的結(jié)果，等等。

　　【浪潮SmartRack中的GPU節(jié)點(diǎn)】

　　浪潮SmartRack協(xié)處理加速整機(jī)柜服務(wù)器實(shí)現(xiàn)了在1U空間里部署4個(gè)Tesla® GPU 加速器，實(shí)現(xiàn)“CPU+協(xié)處理器”協(xié)同計(jì)算加速。此外，該產(chǎn)品還融合了廣泛使用的NVIDIA® CUDA®并行計(jì)算平臺(tái)以及cuDNN GPU加速庫。

　　【科學(xué)計(jì)算到底怎么算】

　　最后，則是本篇壓軸的知識(shí)，那就是，那些個(gè)科學(xué)計(jì)算到底都在算什么東西，又是怎么映射到多線程的？舉個(gè)例子，分子動(dòng)力學(xué)，模擬蛋白質(zhì)折疊過程。（關(guān)于蛋白質(zhì)分子相關(guān)背景知識(shí)請(qǐng)參考筆者的《生物大分子是如何“計(jì)算”的》）一文。一維肽鏈被核糖體讀取DNA加工生成之后，在原子間相互作用力的影響之下，自然折疊卷曲成三維立體結(jié)構(gòu)?？茖W(xué)家們想看看這個(gè)過程到底是怎么運(yùn)動(dòng)的，以及是不是真的可以自行折疊，還是必須依靠其他分子的輔助折疊，于是就想直接根據(jù)原子間作用力算出來。高中物理題，我們隨便計(jì)算個(gè)受力分解，都會(huì)捶胸頓足“這題太難了?。?amp;rdquo;，可想而知，計(jì)算幾百上千個(gè)原子之間的相互作用得有多復(fù)雜。其本質(zhì)上其實(shí)就是牛頓運(yùn)動(dòng)力學(xué)，對(duì)F=ma的求解，一個(gè)原子會(huì)收到各個(gè)方向的各種力，靜電力、化學(xué)鍵、范德華力等等，求出最終的合力方向，則知道該原子即將向哪移動(dòng)，速度為多少。但是隨著所有原子的移動(dòng)，有些快有些慢，各自又朝著不同方向移動(dòng)，那么每個(gè)原子的受力狀況又會(huì)隨之改變，這種問題的求解想想都復(fù)雜，需要使用積分的思想了。積分就是算出每一小步的結(jié)果然后積累成一大步，每一小步范圍內(nèi)可以近似認(rèn)為所有原子的受力狀況暫且沒有變化。這一小步精確到多大呢？目前計(jì)算分子動(dòng)力學(xué)領(lǐng)域一般在飛秒級(jí)，十幾個(gè)飛秒，也就是將合力加在某個(gè)原子上十幾個(gè)飛秒，計(jì)算每個(gè)原子的移動(dòng)坐標(biāo)，這十幾個(gè)飛秒內(nèi)可以近似認(rèn)為原子的受力仍然不變。但是本質(zhì)上，任何微小的變化都會(huì)持續(xù)影響受力，但是大自然底層是如何做到極限精確且連續(xù)的，或者也有某種最小精細(xì)單位，我們無從知曉，或許積分的思想也正預(yù)示著自然底層的確就是有最小移動(dòng)單位的，比如：一個(gè)場。具體可以參考筆者《時(shí)空參悟》一文。模擬完這一小步，然后再次根據(jù)各原子的空間坐標(biāo)，生成新的合力，再走一小步，最終走到某個(gè)穩(wěn)定的點(diǎn)，引力斥力平衡，不能走為止，整個(gè)過程需要龐大的計(jì)算量。

　　那么上述過程如何映射到多個(gè)線程并行運(yùn)算？比如，可以以原子為單位，每個(gè)線程負(fù)責(zé)計(jì)算每個(gè)原子在十幾飛秒后將移動(dòng)到哪個(gè)位置，該線程的輸入值是該原子的元素、化合價(jià)等等，以及初始三維坐標(biāo)和初速度（初始速度為0）；輸出值則是在牛頓力學(xué)公式的作用之下經(jīng)過十幾飛秒加速之后該原子的新三維坐標(biāo)和速度矢量。該線程中又會(huì)有大量函數(shù)相互作用，比如：有函數(shù)會(huì)專門根據(jù)當(dāng)前原子的化合價(jià)以及與其化合的其他原子都是誰，然后計(jì)算靜電力，有的函數(shù)則負(fù)責(zé)計(jì)算氫鍵力、范德華力等。然后求合力F，算出初速度為0，合力F，質(zhì)量m，t=10fs之后的該原子的位置和速度，最后這一步相信高中物理及格的朋友都可以算出來了。這一步結(jié)束之后，數(shù)千個(gè)線程之間便會(huì)發(fā)生通信操作，比如，原子A---原子B----原子C，當(dāng)原子A移動(dòng)位置之后，需要將其位置信息告訴原子B，因?yàn)锳的移位會(huì)對(duì)B的受力造成影響；同理，原子B的最新位置也需要通告給A和C。看模擬的精度，如果認(rèn)為A和C之間的相互作用力不可忽略，那么A和C之間也需要相互通告。此時(shí)可以采用MPI Send操作各自通告結(jié)果，第一輪相互同步完成之后，繼續(xù)再進(jìn)行下一個(gè)10fs的計(jì)算，繼續(xù)循環(huán)下去，直到到達(dá)規(guī)定的時(shí)間或者觸發(fā)條件，則全部線程結(jié)束運(yùn)算，各自執(zhí)行MPI Gather向主線程匯報(bào)最終的計(jì)算結(jié)果，主線程根據(jù)所有原子的三維坐標(biāo)繪圖，最終生成模擬之后的分子空間構(gòu)象。

　　如果用GPU來完成上述計(jì)算，則需要調(diào)用CUDA計(jì)算框架來完成。先把主計(jì)算程序?qū)懞?，也就是從?shù)組中提取其中一個(gè)元素進(jìn)行計(jì)算。然后向GPU分配顯存，再將數(shù)組拷貝到顯存，然后向GPU聲明上述核心函數(shù)需要用多少個(gè)線程來計(jì)算，每個(gè)線程算哪些數(shù)組元素，算完之后將結(jié)果拷貝回主存，形成新的數(shù)組，然后再進(jìn)行一輪。或者直接在GPU內(nèi)部實(shí)現(xiàn)各個(gè)線程之間的數(shù)據(jù)同步，這需要利用到GPU內(nèi)的Share Memory，細(xì)節(jié)筆者不再介紹，有興趣自行了解。

　　分子動(dòng)力學(xué)其實(shí)是最容易理解的超算過程，上述過程也是一種簡化描述，實(shí)際上非常復(fù)雜。有些更加復(fù)雜的科學(xué)計(jì)算，比如冷凍電鏡三維構(gòu)象重構(gòu)。其產(chǎn)生的背景是，分子生物界的科學(xué)家們有時(shí)候不相信用計(jì)算機(jī)模擬計(jì)算出來的分子構(gòu)象，只相信肉眼所見的實(shí)際結(jié)構(gòu)。于是人們想了一種辦法，把含有大量某蛋白質(zhì)分子的液體液氮冷卻成冰塊，然后切割成兩半，用電子顯微鏡對(duì)橫截面拍照，在很大幾率上便會(huì)拍到位于各種角度的蛋白質(zhì)大分子，有豎著的、橫著的、躺著的、各種角度的。然后就將這幅高清圖片輸入到計(jì)算機(jī)進(jìn)行處理，先將這些輪廓進(jìn)行采樣和描述，然后開始用這些各種角度的樣子拼出一副三維空間構(gòu)象。至于程序是采用什么算法將二維輪廓重構(gòu)成三維構(gòu)象的，其計(jì)算過程已經(jīng)超出了筆者的理解范圍，筆者的腦子只能理解F=ma，而這個(gè)過程則需要傅里葉變換，將平面圖變成三維描述，這已經(jīng)完全超越了筆者的認(rèn)知，有興趣可以自行研究，并且分享一下，不亦樂乎？

　　如果你耐心看到這里，證明你的收益達(dá)到了筆者的預(yù)期，那么是不是可以賞了？點(diǎn)擊結(jié)尾贊賞按鈕打賞筆者。