Silicon-Crystal應(yīng)用的神威OpenACC移植與數(shù)據(jù)流驅(qū)動(dòng)任務(wù)圖并行化

(山東科技大學(xué) 計(jì)算機(jī)科學(xué)與工程學(xué)院，山東 青島 266590)

分子動(dòng)力學(xué)(molecular dynamics，MD)模擬是指使用數(shù)值方法，利用計(jì)算機(jī)模擬原子核和電子所構(gòu)成的多體系統(tǒng)的運(yùn)動(dòng)過(guò)程，已被廣泛應(yīng)用于物理、化學(xué)、生物、材料、醫(yī)學(xué)等多個(gè)領(lǐng)域，用來(lái)研究系統(tǒng)的結(jié)構(gòu)和性質(zhì)[1]。在材料領(lǐng)域虛擬過(guò)程工程中,Silicon-Crystal應(yīng)用是研究硅晶體熱傳導(dǎo)性的MD模擬應(yīng)用，有限的計(jì)算能力一直是制約模擬效率的瓶頸[2]。近年來(lái)，高性能計(jì)算技術(shù)的發(fā)展為材料領(lǐng)域的虛擬過(guò)程工程提供了可能[3]。

“神威·太湖之光”是世界上首臺(tái)運(yùn)行速度超過(guò)十億億次的超級(jí)計(jì)算機(jī)，也是我國(guó)第一臺(tái)全部采用國(guó)產(chǎn)處理器SW26010構(gòu)建的超級(jí)計(jì)算機(jī)[4]。清華大學(xué)付昊桓等[5]在“神威·太湖之光”上，利用OpenACC移植大氣模型CAM應(yīng)用，單核組內(nèi)實(shí)現(xiàn)2倍加速，但未進(jìn)行移植后優(yōu)化；上海交通大學(xué)王一超等[6]利用OpenACC移植并優(yōu)化了磁約束聚變領(lǐng)域GTC-P應(yīng)用，單核組內(nèi)實(shí)現(xiàn)2.5倍加速，但缺少對(duì)訪存密集型應(yīng)用帶寬訪存優(yōu)化；中國(guó)科學(xué)院計(jì)算應(yīng)用研究中心張帥等[7]在GPU平臺(tái)上對(duì)MD模擬進(jìn)行訪存優(yōu)化，但未提供模擬中跨時(shí)間迭代問(wèn)題的解決方法。CPU、GPU架構(gòu)與SW26010架構(gòu)存在著差異。SW26010采用片上計(jì)算陣列群和分布式共享存儲(chǔ)相結(jié)合的異構(gòu)眾核體系架構(gòu)，使得MD模擬應(yīng)用的移植具有更大的靈活性，但也使得移植難度加大，目前對(duì)MD模擬移植到“神威·太湖之光”超級(jí)計(jì)算機(jī)上的相關(guān)研究尚未見(jiàn)到。

本文設(shè)計(jì)了一種SW26010主從計(jì)算并行化方案，實(shí)現(xiàn)對(duì)Silicon-Crystal應(yīng)用的神威OpenACC移植與優(yōu)化；以數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行化方法解決任務(wù)間的峰值訪存、跨時(shí)間迭代問(wèn)題，針對(duì)該應(yīng)用訪存密集型特點(diǎn)進(jìn)行帶寬訪存優(yōu)化。

1 背景介紹

1.1 SW26010處理器架構(gòu)及神威OpenACC執(zhí)行模型

“神威·太湖之光”是中國(guó)自主研發(fā)的超級(jí)計(jì)算機(jī)，峰值性能為125.4 PFlops，實(shí)測(cè)峰值約為93 PFlops。采用新一代的眾核異構(gòu)處理器SW26010(架構(gòu)如圖1)。神威OpenACC程序的執(zhí)行模型是在主核指導(dǎo)下，主從核協(xié)同工作，其加速執(zhí)行模型如圖2所示。

圖1 “SW26010”異構(gòu)眾核架構(gòu)Fig.1 Heterogeneous multi-core processor architecture of “SW26010”

SW26010異構(gòu)眾核架構(gòu)中，各核組之間采用片上網(wǎng)絡(luò)互連，每個(gè)核組包含1個(gè)主核(management processing element MPE)、1個(gè)從核簇(8×8=64個(gè)，computing processing element，CPE)、1個(gè)協(xié)議處理單元和1個(gè)內(nèi)存控制器。核組內(nèi)采用共享存儲(chǔ)架構(gòu)，內(nèi)存與主、從核之間可通過(guò)內(nèi)存控制器傳輸數(shù)據(jù)，處理器可通過(guò)系統(tǒng)接口與外部設(shè)備相連[8]。申威眾核處理器旨在用少量具備指令級(jí)并行能力的管理核心集成眾多面向計(jì)算開(kāi)發(fā)的精簡(jiǎn)運(yùn)算核心高效處理線程級(jí)并行，從而大幅提高芯片性能[9]。

程序首先在MPE上啟動(dòng)，以一個(gè)主線程串行執(zhí)行，計(jì)算密集區(qū)域則在主線程的控制下作為加速任務(wù)被加載到加速設(shè)備CPE上執(zhí)行[10]。任務(wù)的執(zhí)行過(guò)程包括：在CPE設(shè)備內(nèi)存上分配所需的數(shù)據(jù)空間；加載任務(wù)代碼至CPE；任務(wù)將所需的數(shù)據(jù)從MPE傳輸至CPE內(nèi)存；等待數(shù)據(jù)傳輸完成；CPE進(jìn)行計(jì)算并將計(jì)算結(jié)果傳送回主存；釋放設(shè)備上的數(shù)據(jù)空間等步驟。

圖2 神威OpenACC執(zhí)行模型Fig.2 Execution model of the Sunway OpenACC

MPE加載一系列任務(wù)到加速設(shè)備上同時(shí)執(zhí)行，但這種fork-join模式在訪存帶寬有限的SW26010處理器上易產(chǎn)生峰值訪存問(wèn)題，使CPE之間相互爭(zhēng)搶帶寬，從而影響計(jì)算性能。

1.2 AceMesh編程框架

AceMesh編程框架是面向網(wǎng)格應(yīng)用[11-12]、以數(shù)據(jù)為中心，應(yīng)用于多核、眾核平臺(tái)上的數(shù)據(jù)流驅(qū)動(dòng)并行編程框架。AceMesh并行編程框架通過(guò)底層的任務(wù)調(diào)度系統(tǒng)[13](運(yùn)行時(shí)庫(kù))對(duì)網(wǎng)格應(yīng)用進(jìn)行任務(wù)圖并行，其核心設(shè)計(jì)思想來(lái)源于圖論中的有向無(wú)環(huán)圖(directed acyclic graph，DAG)。

任務(wù)調(diào)度系統(tǒng)采用探測(cè)-執(zhí)行(inspector-executor)兩階段執(zhí)行的并行模式[14]，該模式對(duì)并行區(qū)域進(jìn)行代碼級(jí)調(diào)度。探測(cè)階段將代碼區(qū)域的控制流和數(shù)據(jù)流信息提交給運(yùn)行時(shí)系統(tǒng)，由運(yùn)行時(shí)系統(tǒng)根據(jù)任務(wù)間依賴關(guān)系建立任務(wù)依賴圖。執(zhí)行階段以構(gòu)建的任務(wù)圖為基礎(chǔ)，依據(jù)資源配置及利用率搭配不同的任務(wù)調(diào)度策略和算法，動(dòng)態(tài)的調(diào)度并行任務(wù)。

圖3 AceMesh編程框架任務(wù)調(diào)度系統(tǒng)Fig.3 Task scheduling system ofAceMesh programming framework

AceMesh編程框架任務(wù)調(diào)度系統(tǒng)結(jié)構(gòu)圖如圖3所示，該調(diào)度系統(tǒng)包括四層：

1) 用戶接口層，收集任務(wù)粒度[15]的描述、數(shù)據(jù)流信息、任務(wù)構(gòu)造等信息；

2) 任務(wù)構(gòu)建層，根據(jù)上層用戶提供的信息，在系統(tǒng)內(nèi)部產(chǎn)生任務(wù)、建立依賴和進(jìn)行任務(wù)圖管理；

3) 任務(wù)調(diào)度層，通過(guò)靜態(tài)調(diào)度和動(dòng)靜結(jié)合調(diào)度兩種方式提供任務(wù)調(diào)度支持。靜態(tài)調(diào)度采用輪詢法按權(quán)值將任務(wù)分配至線程；動(dòng)靜結(jié)合調(diào)度指靜態(tài)調(diào)度策略與任務(wù)竊取調(diào)度算法[13，16]相結(jié)合，提高任務(wù)數(shù)據(jù)重用率和線程間負(fù)載均衡性；

4) 隊(duì)列調(diào)度層，利用線程庫(kù)對(duì)線程私有并發(fā)任務(wù)調(diào)度隊(duì)列進(jìn)行任務(wù)級(jí)調(diào)度。

2 Silicon-Crystal應(yīng)用分析及移植方案設(shè)計(jì)

2.1 應(yīng)用算法和模擬過(guò)程數(shù)據(jù)特征分析

MD模擬中，通過(guò)差分求解牛頓運(yùn)動(dòng)方程可得到系統(tǒng)中原子的一系列位形。由于模擬過(guò)程中力的計(jì)算工作量很大，常用的龍格-庫(kù)塔法已不再適用，Silicon-Crystal應(yīng)用中的TP(Tersoff Potent)模塊利用leap-frog算法[17]模擬原子在Tersoff勢(shì)能作用下的運(yùn)動(dòng)軌跡，在所有的線性微分方程的求解器中都有應(yīng)用。

基于有限差分法leap-frog算法，求解線性常微分方程公式如下：

(1)

(2)

其中，r、V、m、F分別為原子的位置矢量、速度、質(zhì)量、所受勢(shì)能力，Δt為計(jì)算時(shí)間步長(zhǎng)。

2.2 TP模塊并行方案

加速線程庫(kù)(athread庫(kù))是針對(duì)主從加速編程模型所設(shè)計(jì)的程序加速庫(kù)，旨在使用戶能夠方便、快捷地使用核組內(nèi)的線程進(jìn)行控制和調(diào)度，從而更好地發(fā)揮組內(nèi)多計(jì)算核的性能。本研究使用加速線程庫(kù)將TP模塊移植到從核的運(yùn)算模式如圖4所示。

圖4 TP模塊并行方案設(shè)計(jì)Fig.4 Parallel scheme design of TP module

TP模塊的移植主要分為：

1) 計(jì)算網(wǎng)絡(luò)劃分。MPE端沿三維空間x、y和z三個(gè)方向?qū)?shù)據(jù)區(qū)域按比例分成若干矩形體，每一矩形體計(jì)算視為一個(gè)任務(wù)。這樣的劃分方式有兩個(gè)好處：其一，CPE端得到的數(shù)據(jù)在空間上是連續(xù)的，數(shù)據(jù)塊訪問(wèn)開(kāi)銷比較小；其二，分塊內(nèi)中心原子占比相對(duì)較高，減少分塊間的原子通信量，提升計(jì)算效率。

2) 初始化環(huán)境。CPE端對(duì)劃分后任務(wù)內(nèi)的原子信息進(jìn)行初始化，初始化信息包括原子的位置矢量、加速度、速度等。

3) 計(jì)算參數(shù)初始化。初始化MPE端對(duì)Tersoff勢(shì)能下的離散計(jì)算參數(shù)。

4) 力場(chǎng)計(jì)算。以任務(wù)為基本單位將原子信息加載至CPE端進(jìn)行加速計(jì)算，首先進(jìn)行MPE端至CPE端的數(shù)據(jù)拷貝，其次利用CPE端的計(jì)算陣列群加速核心計(jì)算，最后將計(jì)算后的各個(gè)任務(wù)原子信息由CPE端傳回MPE端。

5) 同步力場(chǎng)數(shù)據(jù)。MPE端按照鄰居關(guān)系索引表進(jìn)行任務(wù)間數(shù)據(jù)更新操作，保證數(shù)據(jù)全局一致性。

6) 更新殘量和輸出文件信息。CPE端進(jìn)行每個(gè)時(shí)間步計(jì)算后的殘量更新，MPE端將計(jì)算范數(shù)值輸出至文件系統(tǒng)。

3 神威OpenACC從核移植

神威OpenACC并行編程模型，用編譯指示的方式把應(yīng)用中可并行化的計(jì)算循環(huán)移植到申威處理器從核以加速計(jì)算。具體到Silicon-Crystal應(yīng)用的從核移植，主要分為以下三個(gè)步驟：

1) 循環(huán)并行化。Silicon-Crystal應(yīng)用以任務(wù)分片存儲(chǔ)的數(shù)據(jù)為基本單位進(jìn)行模擬計(jì)算，在分塊級(jí)的for循環(huán)上添加相應(yīng)的指導(dǎo)語(yǔ)句#pragma acc parallel loop，將計(jì)算部署在64個(gè)從核上并行執(zhí)行。

需要注意的是，gang、worker、vecotr是OpenACC2.0中的3層循環(huán)設(shè)計(jì)，由于神威眾核架構(gòu)在物理上并沒(méi)有分層需求，所以神威OpenACC的實(shí)現(xiàn)是把gang設(shè)置成64，worker設(shè)為1。

2) 基于計(jì)算數(shù)據(jù)優(yōu)先的數(shù)據(jù)管理。神威眾核架構(gòu)中存在訪存帶寬較小的問(wèn)題，故從核移植并行化過(guò)程最為關(guān)鍵的是將加速計(jì)算需要的數(shù)據(jù)提前拷貝到訪問(wèn)延遲低的SPM(scratch pad memory)。本研究采用計(jì)算數(shù)據(jù)優(yōu)先傳輸策略即將所有計(jì)算涉及的數(shù)據(jù)優(yōu)先傳至SPM。

計(jì)算數(shù)據(jù)優(yōu)先傳輸過(guò)程如下：

i) 按循環(huán)索引劃分傳輸。若數(shù)組的索引變量與循環(huán)索引變量緊耦合時(shí)，神威OpenACC編譯器將數(shù)組劃分為64份，然后利用DMA的方式將劃分后的數(shù)據(jù)集中傳輸至各從核SPM中，并將任務(wù)內(nèi)的鄰居關(guān)系表、打包后的計(jì)算參數(shù)順序傳遞給從核。具體使用copy/copyin/cpoyout等指導(dǎo)語(yǔ)句完成(如圖5)。

①#pragma acc parallel loop?②local(pi) //變量局存私有化③copy(ax,ay,az,pox,poy,poz)annotate(dimension(ax(PN,N),ay(PN,N),az(PN,N))) //按循環(huán)索引傳輸④copyin(maplis) annotate(dimension(maplist(4?N))) //鄰居關(guān)系表傳輸⑤packin(deltaT,soma,Mu,Lamda,Beta,nexp,InteriorCutoff,ExteriorCutoff) //離散計(jì)算參數(shù)打包傳輸⑥for (pi = 0; pi < PN ; pi++){…}

圖5 數(shù)據(jù)管理過(guò)程的函數(shù)指導(dǎo)語(yǔ)言實(shí)現(xiàn)
Fig.5 Implementation of functional instruction language based on data management

ii) 變量局存私有化。對(duì)于并行循環(huán)索引變量等線程私有變量，既可使用private子句也可使用local子句將變量私有化，考慮到private是線程私有化變量，變量值仍在主存中，而local是線程私有化的局存變量，存儲(chǔ)在SPM中，數(shù)據(jù)訪問(wèn)更加高效，故采用local子句進(jìn)行變量的私有化。

iii) 離散計(jì)算參數(shù)打包傳輸。Silicon-Crystal應(yīng)用存在多個(gè)離散標(biāo)量的模擬參數(shù)需要傳送至從核，若一一傳輸需要頻繁的使用DMA方式，會(huì)大大增加訪存開(kāi)銷。在此情況下，本文利用pack/packin/packout等神威定制的指導(dǎo)語(yǔ)句將離散數(shù)據(jù)打包后一次傳遞，以更充分有效地利用有限的訪存帶寬。

綜上，得到 Silion-Crystal應(yīng)用移植中的數(shù)據(jù)管理過(guò)程的函數(shù)指導(dǎo)語(yǔ)言實(shí)現(xiàn)如圖5。

3) 加速代碼區(qū)約束處理。SWACC編譯器進(jìn)行OpenACC并行化過(guò)程中，對(duì)并行區(qū)的代碼有一定的要求。如在加速區(qū)代碼中存在函數(shù)調(diào)用時(shí)，需在函數(shù)定義處添加routine子句指示，否則生成從核代碼將找不到函數(shù)的位置。但目前routine子句只適用Fortran程序，C代碼暫不支持。Silicon-Crystal應(yīng)用程序是C代碼程序，無(wú)法利用routine子句修飾從核函數(shù)。本研究通過(guò)利用宏定義實(shí)現(xiàn)力場(chǎng)計(jì)算的內(nèi)聯(lián)函數(shù)，來(lái)解決移植過(guò)程中加速區(qū)函數(shù)返回值異常的問(wèn)題。

圖6 神威OpenACC移植性能數(shù)據(jù)Fig.6 Performance data of the Sunway OpenACC transplant

將運(yùn)行在1個(gè)主核上的Silicon-Crystal作為測(cè)試基準(zhǔn)，分別與循環(huán)并行化、基于計(jì)算數(shù)據(jù)優(yōu)先的訪存和離散計(jì)算參數(shù)打包傳輸3個(gè)方面在單核組上進(jìn)行性能測(cè)試(圖6)。測(cè)試問(wèn)題規(guī)模：回環(huán)中存在131 072個(gè)粒子，迭代計(jì)算次數(shù)為1 000次。

可以看出，對(duì)于訪存密集型的應(yīng)用，僅進(jìn)行循環(huán)并行化將計(jì)算過(guò)程移至CPE端，性能反而會(huì)降低；按照計(jì)算數(shù)據(jù)優(yōu)先方式通過(guò)DMA方式放入從核SPM中，性能開(kāi)始超越主核；通過(guò)pack子句對(duì)離散計(jì)算參數(shù)打包后再傳輸，性能進(jìn)一步提升，整體應(yīng)用較主核版實(shí)現(xiàn)了2.26倍的加速。

4 數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行化

AceMesh任務(wù)調(diào)度系統(tǒng)的設(shè)計(jì)思想來(lái)源于數(shù)據(jù)結(jié)構(gòu)中的有向無(wú)環(huán)圖，即任務(wù)依賴圖。任務(wù)依賴圖在圖論中是指：如果一個(gè)有向圖無(wú)法從某個(gè)頂點(diǎn)出發(fā)經(jīng)過(guò)若干條邊回到該點(diǎn)，則這個(gè)圖是一個(gè)任務(wù)依賴圖。任務(wù)依賴圖中的頂點(diǎn)代表任務(wù)，圖中的邊代表任務(wù)間的依賴關(guān)系。根據(jù)任務(wù)依賴圖的特點(diǎn)，將并行計(jì)算中的大規(guī)模計(jì)算問(wèn)題劃分為N(N≥1)個(gè)任務(wù)，并根據(jù)各個(gè)任務(wù)的依賴關(guān)系建立任務(wù)依賴圖，圖中所有沒(méi)有后繼的頂點(diǎn)都執(zhí)行完后，任務(wù)依賴圖的執(zhí)行完成。

在神威眾核處理器上任務(wù)圖并行化過(guò)程分為任務(wù)構(gòu)圖期和任務(wù)執(zhí)行期。構(gòu)圖期是任務(wù)構(gòu)建的探測(cè)過(guò)程，旨在根據(jù)注冊(cè)的數(shù)據(jù)地址去建立任務(wù)間的依賴關(guān)系，在不改變串行序結(jié)果的情況下以數(shù)據(jù)流調(diào)整執(zhí)行序列；執(zhí)行期是指按照構(gòu)圖期間構(gòu)建的DAG圖，搭配任務(wù)調(diào)度系統(tǒng)的不同調(diào)度策略執(zhí)行任務(wù)的過(guò)程。故任務(wù)圖并行化總時(shí)間等于構(gòu)圖時(shí)間(graph time)加執(zhí)行時(shí)間(execution time)。

Silicon-Crystal應(yīng)用在太湖之光上使用任務(wù)圖并行化主要分為以下3個(gè)步驟：

1)主核構(gòu)建任務(wù)依賴圖。根據(jù)不同并行區(qū)內(nèi)劃分的任務(wù)按照對(duì)內(nèi)部訪問(wèn)的數(shù)據(jù)依賴關(guān)系進(jìn)行地址注冊(cè)，構(gòu)建出任務(wù)執(zhí)行序DAG圖。

圖7為Silicon-Crystal應(yīng)用2線程4任務(wù)依賴圖。其中，每個(gè)橢圓代表包裝后的一個(gè)任務(wù)，橢圓中第一個(gè)數(shù)字為并行區(qū)編號(hào)，第二個(gè)數(shù)字為任務(wù)編號(hào)；箭頭代表任務(wù)間的數(shù)據(jù)訪問(wèn)先后的依賴關(guān)系；陰影、非陰影圓圈代表執(zhí)行時(shí)不同的線程；實(shí)線邊是任務(wù)垂直后繼依賴邊，虛線邊是普通后繼依賴邊。在任務(wù)執(zhí)行期間，采用的調(diào)度策略使垂直后繼任務(wù)優(yōu)先于普通后繼任務(wù)執(zhí)行，旨在使任務(wù)間的數(shù)據(jù)重用得到最大化。此外，截?cái)嗖⑿袇^(qū)間任務(wù)執(zhí)行的依賴關(guān)系，按照與神威OpenACC相同的控制流驅(qū)動(dòng)的fork-join執(zhí)行模式，稱為任務(wù)圖單步執(zhí)行。

圖7 2線程4任務(wù)時(shí)TP模塊任務(wù)依賴圖Fig.7 Task dependency graph of TP module under 2 threads 4 tasks

圖8 神威OpenACC與任務(wù)圖并行化性能Fig.8 Performance of the Sunway Open ACC andtask graph parallelization

2)從核包裝任務(wù)函數(shù)，將應(yīng)用主要的計(jì)算代碼包裝成任務(wù)函數(shù)。從核任務(wù)函數(shù)根據(jù)構(gòu)圖期分配的函數(shù)參數(shù)、循環(huán)劃分尺寸、數(shù)據(jù)區(qū)劃分尺寸等信息，包裝任務(wù)函數(shù)，放入從核陣列并行計(jì)算。

3)從核數(shù)據(jù)管理。SW26010主從核間的數(shù)據(jù)傳輸通過(guò)DMA實(shí)現(xiàn)，DMA只能由從核發(fā)起，主核被動(dòng)進(jìn)行數(shù)據(jù)傳輸。從核制定傳輸?shù)哪Ｊ綍r(shí)，數(shù)據(jù)傳輸依據(jù)數(shù)據(jù)在主存數(shù)據(jù)區(qū)內(nèi)存儲(chǔ)地址的連續(xù)性和從核計(jì)算實(shí)際需要的數(shù)據(jù)尺寸進(jìn)行傳輸。DMA數(shù)據(jù)傳輸方式分為跨步式數(shù)據(jù)傳輸和非跨步式數(shù)據(jù)傳輸。兩種傳輸模式下，軟件開(kāi)銷主要體現(xiàn)在傳輸?shù)膯?dòng)和對(duì)DMA傳輸回答字的處理。本研究采用數(shù)據(jù)分片存儲(chǔ)的數(shù)據(jù)結(jié)構(gòu)，將任務(wù)訪問(wèn)的數(shù)據(jù)進(jìn)行分片存儲(chǔ)，并以數(shù)據(jù)塊編號(hào)為索引劃定數(shù)據(jù)區(qū)，通過(guò)athread_get/athread_put接口進(jìn)行非跨步傳輸；對(duì)于離散數(shù)據(jù)訪問(wèn)，在主核代碼中對(duì)離散數(shù)據(jù)打包后，使用DMA方式進(jìn)行數(shù)據(jù)傳輸，最后在從核代碼中對(duì)數(shù)據(jù)進(jìn)行解包，提高離散數(shù)據(jù)的訪問(wèn)效率。

本節(jié)將第3節(jié)中OpenACC優(yōu)化后的版本(ACC)作為基礎(chǔ)版，同等的優(yōu)化條件下，與任務(wù)圖單步版(single step of DAG)、任務(wù)圖亂序版(Unordered DAG)進(jìn)行實(shí)驗(yàn)對(duì)比。迭代時(shí)間步長(zhǎng)為1時(shí)，Silicon-Crystal應(yīng)用性能如圖8所示。

可以看出，由于任務(wù)圖單步版和OpenACC采用相同的fork-join模式，即并行區(qū)開(kāi)始spawn線程，并行區(qū)結(jié)束wait所有線程，故二者的執(zhí)行時(shí)間一致，說(shuō)明兩者具有計(jì)算一致性。任務(wù)圖亂序版比任務(wù)單步版提升27%，驗(yàn)證了任務(wù)間的亂序執(zhí)行，可以錯(cuò)開(kāi)峰值帶寬競(jìng)爭(zhēng)，充分的利用從核訪存帶寬。但是，由于神威OpenACC采用fork-join模式，其執(zhí)行時(shí)間即為總時(shí)間；數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行需要在構(gòu)圖期構(gòu)建任務(wù)執(zhí)行的依賴關(guān)系，故其總時(shí)間為構(gòu)圖時(shí)間與執(zhí)行時(shí)間之和。實(shí)驗(yàn)結(jié)果表明，Silicon-Crystal應(yīng)用的任務(wù)圖并行化存在相對(duì)總時(shí)間8%的構(gòu)圖時(shí)間，加上此部分構(gòu)圖開(kāi)銷，總時(shí)間上任務(wù)圖并行比ACC性能提升11.5%。

表1 多時(shí)間步擴(kuò)展下任務(wù)圖并行化構(gòu)時(shí)間Tab.1 Times of task graph parallelization based on multiple time steps

圖9 多時(shí)間步下任務(wù)圖并行化加速比Fig.9 Acceleration ratio of task graph parallelizationunder multiple time steps

傳統(tǒng)的fork-join模式無(wú)法擴(kuò)展多時(shí)間步的迭代計(jì)算，任務(wù)圖卻可打通迭代時(shí)間步間的并行區(qū)域，即在構(gòu)圖期依據(jù)多個(gè)時(shí)間步下任務(wù)的數(shù)據(jù)流構(gòu)建出任務(wù)依賴圖，執(zhí)行期按多時(shí)間步下任務(wù)亂序調(diào)度方式執(zhí)行任務(wù)，從而將迭代時(shí)間步由單時(shí)間步擴(kuò)展至多時(shí)間步，如表1所示，隨著任務(wù)圖并行在多時(shí)間步的擴(kuò)展，執(zhí)行時(shí)間進(jìn)一步降低，構(gòu)圖時(shí)間逐漸降低，使得任務(wù)圖并行性能進(jìn)一步提升。

多時(shí)間步下任務(wù)圖并行化加速比如圖9所示，以主核版作為基準(zhǔn)版，使用神威OpenACC移植利用從核加速，實(shí)現(xiàn)2.26倍加速比；時(shí)間步為1時(shí)任務(wù)圖并行加速比為2.52；隨著時(shí)間步的擴(kuò)展，任務(wù)圖規(guī)模隨之增加，任務(wù)的亂序使錯(cuò)峰訪存的優(yōu)勢(shì)進(jìn)一步擴(kuò)大，時(shí)間步擴(kuò)展至20時(shí)趨于平穩(wěn)，加速比達(dá)到3.2。

5 總結(jié)與未來(lái)工作

本研究為Silicon-Crystal應(yīng)用設(shè)計(jì)了一套在SW26010上實(shí)現(xiàn)主從計(jì)算的并行化方案，利用OpenACC完成了向目標(biāo)平臺(tái)“神威·太湖之光”上的移植，在單核組內(nèi)實(shí)現(xiàn)了2.52倍加速；針對(duì)該應(yīng)用訪存密集的行為特點(diǎn)，以數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行化方法解決任務(wù)間的峰值訪存和跨時(shí)間迭代問(wèn)題，結(jié)果表明，Silicon-Crystal應(yīng)用在數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行在單時(shí)間步下性能提升11.5%，多時(shí)間步下性能提升42%，總體較主核實(shí)現(xiàn)3.2倍加速。

數(shù)據(jù)流驅(qū)動(dòng)的任務(wù)圖并行編程模型采用AceMesh任務(wù)調(diào)度系統(tǒng)中的低級(jí)接口對(duì)程序源碼進(jìn)行優(yōu)化，隨著本課題組數(shù)據(jù)驅(qū)動(dòng)的并行調(diào)度系統(tǒng)自動(dòng)轉(zhuǎn)譯器的完善，未來(lái)將使用指導(dǎo)語(yǔ)言的高級(jí)形式對(duì)代碼進(jìn)行自動(dòng)源源變換，從而實(shí)現(xiàn)通過(guò)指導(dǎo)語(yǔ)言方式對(duì)應(yīng)用任務(wù)圖并行的自動(dòng)化過(guò)程。