XeonPhi平臺上基于模板優化的3DGVF場計算加速*

齊 金，李 寬，楊燦群，杜云飛

(1.國防科學技術大學并行與分布處理重點實驗室，湖南 長沙 410073;2.國防科學技術大學計算機學院，湖南 長沙 410073)

XeonPhi平臺上基于模板優化的3DGVF場計算加速*

齊 金1，李 寬2，楊燦群1，杜云飛2

(1.國防科學技術大學并行與分布處理重點實驗室，湖南 長沙 410073;2.國防科學技術大學計算機學院，湖南 長沙 410073)

3D梯度向量流場(3D GVF field)廣泛應用于多種3D圖像分析算法中，其計算需要多次迭代，計算量大，如何提高其計算速度具有重要的研究意義。面向Intel Xeon Phi眾核集成架構，首次進行了3D GVF場計算的加速優化。首先，挖掘3D圖像像素點間存在的天然并行性，發揮眾核架構優勢，嘗試線程級并行(多核)和數據級并行(SIMD)。其次，3D GVF場的計算過程是一種典型的3D-7點模板運算，結合Xeon Phi架構的L2 緩存規格，提出一種高效的數據分塊策略，充分挖掘數據的時/空局部性，有效緩解模板計算引起的緩存缺失，提升了計算性能。實驗結果表明，引入模板優化技術能顯著提升3D GVF場的計算速度，在圖像維度為5123時，所提方法在57核Xeon Phi平臺上的性能相比在2.6 GHz 8核16線程的Intel Xeon E5-2670 CPU上的性能，加速比可達2.77。

3D梯度向量流場；Xeon Phi；模板優化；緩存分塊

1 引言

在基于變分法的圖像處理中，主動輪廓模型(Active Contour Model)廣泛應用于邊界檢測、圖像分割和運動跟蹤[1～3]。概括來說，主動輪廓模型將曲線(曲面)與能量函數相聯系，由內力和外力場共同引導，使曲線(曲面)不斷向能量最小化的方向演化。其中，內力由曲線決定，外力場由圖像計算得到。諸多的外力場定義中，Xu C等人[4]提出的梯度向量流場GVF場(Gradient Vector Flow Field)通過一組偏微分方程對圖像的梯度向量進行擴散，具有捕獲范圍大、抗噪聲等優點，成為主動輪廓模型中經典的外力場，得到了廣泛的研究與應用[5,6]。

GVF場的計算與圖像大小規模緊密相關，且需多次迭代。3D圖像數據量大，其3D GVF場的計算速度一直是制約其應用的瓶頸，如何提高GVF場，尤其是3D GVF場的計算速度已獲得研究者的關注。在高性能計算領域，已有研究使用GPU加速3D GVF場計算，如GPU+OpenGL[7]、GPU+OpenCL[8]等。圖像中各像素的處理存在天然的并行性，加上GPU中獨特的紋理存儲(Texture Memory)能方便存取鄰居像素數據，3D GVF場在GPU平臺上取得了較好的加速效果。

Intel 新推出的Xeon Phi眾核集成架構，提供大量IA(Intel Architecture)架構的輕量核，在兼容傳統編程模型的基礎上，能提供更高的計算性能。在此背景下，本文首次面向Xeon Phi平臺進行3D GVF場計算的并行優化。主要的工作包括兩個方面：(1)挖掘3D圖像像素點間存在的天然并行性，發揮眾核架構優勢，嘗試線程級并行(多核)和數據級并行(SIMD)。并對圖像數據在Xeon Phi協處理器內存中的存取模式進行優化，以節省指令、提高吞吐率，這對其他圖像算法在Xeon Phi上的處理均有借鑒意義。(2)計算3D GVF場時，需存取3D空間中的鄰居像素，這是一種典型的3D-7點Laplace模板計算。在Xeon Phi平臺上，結合L2緩存大小，提出一種高效、具體的分塊策略，充分挖掘數據的時/空局部性，有效緩解模板計算引起的緩存缺失，提升了計算性能。實驗結果表明,引入模板優化技術能顯著提升3D GVF場的計算速度。

2 背景和相關工作

(1)3D GVF場計算方法。

3D GVF場是能使得如下能量函數最小的圖像空間的向量場V；

(1)

其中，V0是初始向量場。

上述最優化問題可通過如下歐拉方程解得：

(2)

該問題對應的數值解法如算法1所示：

算法13D GVF數值迭代解法

fori∈[1Iterations] do

forpoint(x,y,z)∈the image do

laplacian←-6Vi(x,y,z)+(x+1,y,z)+Vi(x-1,y,z)+(x,y+1,z)+Vi(x,y-1,z)+(x,y,z+1)+Vi(x,y,z-1)

Vi+1(x,y,z)←Vi(x,y,z)+μ×laplacian-(Vi(x,y,z)-V0(x,y,z))|V0(x,y,z)|2

end for

end for

可見，同次迭代中，各點的梯度向量更新是相對獨立的，這種天然的數據并行性能很好地發揮Xeon Phi的強大計算能力，為后續的線程級并行(多核)和數據級并行(SIMD)奠定了基礎。

(2)模板計算與優化。

所謂模板計算，指的是多次迭代，且在一次迭代內按網格點的順序, 依次對所有網格點進行更新操作, 更新時會用到該網格點的相鄰點的信息。模板計算可按維度和更新所用鄰居點的數目來分類，圖1給出了2D-5點、3D-7點模板示例，結合前面的敘述可知，3D GVF場計算中對每個像素點梯度向量的更新屬于典型的3D-7點模板計算。

Figure 1 Samples of stencil computations圖1 模板計算示例

模板計算的兩個顯著特點是：①不連續的內存訪問模式，容易造成緩存缺失。文獻[9]對模板計算的緩存缺失因素作了詳細的分析，概括來說，當數組大于緩存容量時, 本次更新的數據在下次更新前已經被寫回內存；而且當數據量大時，多次迭代會導致數據緩存的容量缺失(CapacityMiss)。②計算/訪存比低，緩存中數據重用率低，對訪存帶寬要求高。

對模板計算優化的關鍵在于充分開發計算和數據的時/空局部性。諸多優化方法中，緩存分塊(CacheBlocking)是一種典型的優化思路。RiveraG等[10]提出的緩存分塊策略如圖2所示，其中I是單元跨度(Unit-stride)維度，或稱變化最快的維度；K是變化最慢的維度。實驗表明：針對I和J兩個維度的分塊打斷了線程內存讀取流的持續性，不利于數據并行化的展開，因此探索高效的分塊方式，并從XeonPhi體系結構上找到理論支撐，具有重要的應用價值。

Figure 2 Rivera cache blocking圖2 Rivera緩存分塊策略

(3)XeonPhiTMcoprocessors體系結構。

XeonPhi擁有數十個核，每個核包含一個支持512位SIMD的向量處理單元(VPU)，可同時處理8路雙精度或16路單精度浮點數據。Intel提供Intel向量化庫、IntelIntrinsic函數等方式使用該VPU單元。每個核擁有32KB的一級數據緩存和32KB的一級指令緩存。此外，每個核還可以使用512KB的L2級Cache。不同核的L2Cache通過雙向內存控制器相連。

3 GVF并行算法設計

3.1 線程級并行和數據級并行

圖像空間中各像素點計算的天然并行性，使得XeonPhi眾核集成架構的強大計算能力能得到有效發揮。本節討論XeonPhi對圖像處理算法的通用加速方法，主要由兩個層次組成：(1)線程級并行；(2)數據級并行。此外，還對XeonPhi圖像處理中的存取模式進行了有效的探索，可視為XeonPhi對圖像操作的通用預處理。

首先是圖像邊緣點的判斷和處理，由第2節的數值解法可知，計算邊緣點的laplacian時，其鄰居點有的并不存在，為防止內存越界訪問，需要在編碼中引入額外指令判斷當前處理的點是否為邊緣點，勢必會影響整體計算效率。為解決該問題，本文對圖像進行邊界擴充和鏡像填充，向各個方向均擴充一個像素。為方便觀察，圖3以二維圖像的左上角作示例，3D圖像的處理與此一致。

Figure 3 Expanding 2D image with 1 pixel in all directions， the arrows indicates values the new boundary pixels use.圖3 二維圖像邊界擴充與鏡像填充示例， 箭頭指示數據拷貝的方向

為使得XeonPhi的512位VPU可以高效地訪問數據，提高計算吞吐率，采取如下措施：(1)動態分配內存時，使用_mm_malloc函數，確保所分配內存邊界對齊；(2)對單元跨度維度，即變化最快的維度，在圖像邊界擴充的基礎上再次擴展，確保其長度為64字節的整數倍。

根據圖像數據在內存中的組織方式，在變化慢的兩個維度使用線程級并行，使用OPENMP編譯指導語句將任務劃分給多個核；需要說明的是：(1)對XeonPhi而言，線程綁定方式對于計算效率有較大的影響，使用“KMP_AFFINITY=balanced”模式確保所有線程平衡地劃分到Xeon Phi協處理器的核上；(2)使用OPENMP提供的collapse指導語句將兩個維度的循環折疊到一個大的循環中，能有效減少OPENMP任務調度的開銷。

低維度并行方面，使用編譯指導語句確保Xeon Phi 512位的SIMD單元得到有效的利用。

3.2 緩存分塊策略

本節將模板優化相關方法引入3D GVF場計算加速，在綜合考慮3D GVF模板計算特點和Xeon Phi體系結構的基礎上，提出一種Xeon Phi平臺相關的分塊策略。

首先，為充分發揮Xeon Phi中SIMD單元的效能，對單元跨度維度，即變化最快的維度不分塊，如此能保證每個線程的內存讀取流持久而連續；其次，參考文獻[11]，對N×N×N大小的區域，推薦的形狀為(N-2)×s×(s×L/2)，其中，L為一個Cache行的長度，s為分塊大小。對Xeon Phi而言，每個核的L2緩存為512 KB，L=64B，為確保內存讀取各鄰居偏移量項時的局部性，也為確保 Xeon Phi 架構下使用 512 KB L2 高速緩存時的局部性，避免Cache容量缺失，近似有：

N×s×(s×64/2)×Tp×Nm<512KB

(3)

3.3 整體算法(偽代碼)

3D GVF算法加速偽代碼如下所示：

1. for (t=0;t

2. #pragma omp parallel for collapse(2)

3. for(jj=1;jj

4. for(kk=1;kk

5. for(k=kk;k

6. for(j=jj;j

7. #pragma simd

9. …// 3D GVF stencil computation

10. }

11. }

12. }

13. }

14. }

15. }

其中,timesteps指迭代步數;nx、ny、nz分別代表圖像的三個維度大小;s為第二維分塊大小，32s為第三維分塊大小。

1.3 利用果蠅的蛹收集處女蠅 接種純種親本果蠅，當觀察瓶壁上出現較多的黑褐色的蛹時，用干凈解剖針輕輕地把黑褐色的蛹取出，單獨放到10mL的塑料離心管里，置于溫度25℃、濕度60%的培養箱里進行培養。每天觀察，待果蠅羽化出來后進行麻醉，鑒別雌雄，收集處女蠅。

4 實驗結果分析

4.1 實驗平臺

本文以Native模式使用XeonPhi，此時其可視為一個獨立的處理器，完成所有運算，不受CPU控制。此模式下運行的代碼除了512位SIMD指令外，不能含有其他SIMD指令，而且編譯時需要加入編譯選項-mmic。作為對比，本文亦在8核2.60GHz的IntelXeonE5-2670CPU上對3DGVF算法進行了性能測試(采用AVX進行向量化)。XeonPhi與E5-2670CPU的具體配置如表1所示。采用的編譯器為Inteliccversion13.0，OPENMP版本為3.1。

Table 1 System configuration表1 系統配置

為了克服隨機因素的影響，本文所有測試采用執行5遍求平均值的方式進行，每遍執行時迭代100次，針對不同的問題規模，采用關鍵函數墻內時間作為性能度量。

4.2 多線程和SIMD對GVF場計算速度的改進

本小節評測在Xeon Phi上使用多線程和向量化取得的性能提升。圖4是不同線程下，有無向量化時的3D GVF場計算性能對比，其中，橫坐標是所用線程數，縱坐標是墻內計算時間，單位為秒。

Figure 4 Computation performance under different setups圖4 不同配置下的系統計算性能

可以看出，在眾核集成架構中，多線程能顯著提升計算性能，在線程數較低時表現尤其明顯。但是，當線程數達到一定規模時(滿足每核兩個線程時)，線程數目的提升對計算性能的影響不大。對向量化SIMD而言，當線程數較少時，使用向量化能以超過兩倍的加速比提升計算性能。但是，當線程數較多時，SIMD帶來的性能改進已幾乎可忽略不計(<1%)，本次測試中，在線程數較多時，甚至出現了SIMD拖低性能的現象。綜合而言，使用Xeon Phi進行圖像相關算法處理，尤其是像素級stencil遍歷運算時，滿足每核兩個線程是較優的選擇。

4.3 不同分塊

本小節評測提出的分塊策略能否達到較優的效果。與之前諸多采用遍歷不同分塊求最優的工作不同，本文給出了分塊大小的經驗指導公式，如公式(3)所示。為評測該公式的性能，設計如下測試：問題規模選取5123，將第二維分塊的大小即第3.2節中的s從1到8遍歷，第三維分塊的大小選取為n×s，n的取值為1、2、4、8、16、32。本文經驗公式給出的分塊大小為(512,2,64)。

將Xeon Phi的線程數設置為每核兩個線程，總計114個線程。不采用分塊策略時，Xeon Phi計算51233D GVF場所需時間平均值為8.471秒，以此為基準，不同分塊大小相對此基準的加速比如圖5所示，兩個坐標軸分別對應n(橫坐標)和s(縱坐標)的大小。

Figure 5 Speedup ratios under different cache blocking sizes圖5 不同分塊大小相比基準情況的加速比

可以看出，本文所提出的分塊策略相比基準情況有1.25倍的加速比，在所測試的組合中是最優的，從實驗角度驗證了3.2節中對3D GVF計算和Xeon Phi L2緩存的分析。觀察圖5還可發現,加速比相對較高的組合集中在圖5所示矩陣的反對角線方向。

4.4 Xeon Phi與Xeon CPU優化后的計算性能對比

為直觀體現本文在Xeon Phi上對3D GVF場計算的加速效果，本小節將Xeon CPU和Xeon Phi的計算性能進行了比對，兩者所用優化措施均包括多線程和向量化，在Xeon Phi端使用本文提出的經驗分塊公式，在Xeon CPU端以遍歷尋優的方式取最佳性能。兩者的對比結果如表2所示，可見本文的優化方法使得Xeon Phi的絕對性能達到了Xeon CPU的2.77倍。

Table 2 Performance comparisonsbetween Xeon CPU and Xeon Phi表2 Xeon CPU和Xeon Phi的性能比較

5 結束語

本文首次針對Xeon Phi平臺進行了GVF場計算加速研究。GVF場的計算體現出圖像處理中各像素的天然并行性，在合理安排圖像數據內存存取模式的基礎上，結合通用的線程級并行和數據級并行對GVF場進行加速優化；同時，GVF場計算是一種典型的模板計算，結合Xeon Phi二級Cache結構特點，提出了高效的分塊大小經驗公式，避免了費時費力的分塊尋優。實驗結果表明，本文所提方法在Xeon Phi平臺計算3D GVF場取得了很好的加速比。

[1] Kass M, Withkin A, Terzopoulos D. Snakes:Active contour models[J]. International Journal of Computer Vision, 1988,1(4):321-331.

[2] Zhong Y, Jain A K, Dubuisson-Jolly M P. Object tracking using deformable templates[J]. IEEE Transactions on Pattern Analysis and Machine Intelligence, 2000, 22(5):544-549.

[3] Caselles V, Morel J M, Sbert C. An axiomatic approach to image interpolation[J]. IEEE Transactions on Image Processing, 1998, 7(3):376-386.

[4] Xu C, Prince J L. Snakes, shapes, and gradient vector flow[J]. IEEE Transactions on Image Processing, 1998, 7(3):359-369.

[5] Jifeng N,Chengke W,Shigang L,et al.NGVF:An improved external force field for active contour model[J]. Pattern Recognition Letters, 2007, 28(1):58-63.

[6] Wang Y Q, Jia Y D. A novel approach for segmentation of cardiac magnetic resonance images[J]. Chinese Journal of Computers, 2007, 30(1):129-136.(in Chinese)

[7] He Z, Kuester F. GPU-based active contour segmentation using gradient vector flow[M]∥Advances in Visual Computing, Berlin:Springer, 2006:191-201.

[8] Smistad E, Elster A C, Lindseth F. Real-time gradient vector flow on GPUs using OpenCL[J]. Journal of Real-Time Image Processing, 2012,DOI 10.1007/S11554-012-0257-6.

[9] Leopold C. Cache miss analysis of 2D stencil codes with tiled time loop[J]. International Journal of Foundations of Computer Science, 2003, 14(1):39-58.

[10] Rivera G, Tseng C W. Tiling optimizations for 3D scientific computations[C]∥Proc of ACM/IEEE 2000 Conference on Supercomputing, 2000:32.

[11] Leopold C. Tight bounds on capacity misses for 3D stencil codes[C]∥Proc of the International Conference on Computational Science-Part I, 2002:843-852.

附中文參考文獻：

[6] 王元全, 賈云得. 一種新的心臟核磁共振圖像分割方法[J]. 計算機學報, 2007, 30(1):129-136.

QIJin,born in 1988,MS candidate,his research interest includes system software.

李寬(1984-),男，山東寧陽人，博士，助理研究員，研究方向為并行計算和圖像處理。E-mail:likuan@nudt.edu.cn

LIKuan,born in 1984,PhD,assistant researcher,his research interests include parallel computing, and image processing.

楊燦群(1968-),男，湖南桃江人，博士，研究員，研究方向為系統軟件。E-mail:canqun@nudt.edu.cn

YANGCan-qun,born in 1968,PhD,research fellow,his research interest includes system software.

杜云飛(1980-),男，安徽阜南人，博士，助理研究員，研究方向為并行計算、編譯技術和程序性能優化。E-mail:forest80@163.com

DUYun-fei,born in 1980,PhD,assistant researcher,his research interests include parallel computing,compiler technology, and program performance optimization.

Accelerating3DGVFfieldcomputationonXeonPhiusingstenciloptimization

QI Jin1,LI Kuan2,YANG Can-qun1,DU Yun-fei2

(1.National Laboratory of Parallel and Distributed Processing,National University of Defense Technology,Changsha 410073;(2.College of Computer Science,National University of Defense Technology,Changsha 410073,China)

3D Gradient Vector Flow (GVF) field has wide applications in many image processing algorithms. The computation of GVF field typically needs several iterations and is rather time consuming. Therefore, it is important and meaningful to improve the computation speed of 3D GVF field. The data level parallelism and thread level parallelism are introduced to accelerate the GVF field computation procedure on Intel Xeon Phi many core integrated platform for the first time. Meanwhile, GVF field computation is a kind of stencil computation, whose computation-memory access ratio is low. A novel cache blocking strategy is proposed to fully utilize the L2 cache of Xeon Phi architecture，and to improve the computation speed of GVF field. The experimental results show that the proposed optimizations could effectively improve the speed of GVF filed computation. Especially, for a 51233D image, compared with the performance obtained by a 2.6G Hz 8 core 16threads Intel Xeon E5-2670 CPU, the speedup achieved on Xeon Phi is 2.77X.

3D GVF field;Xeon Phi;stencil optimization;cache blocking

1007-130X(2014)08-1435-06

2013-08-12;

：2013-11-11

國家863計劃資助項目(2012AA010903)；國家自然科學基金資助項目(61170049，61303189)

TP393

：A

10.3969/j.issn.1007-130X.2014.08.003

齊金(1988-),男，湖南株洲人，碩士生，研究方向為系統軟件。E-mail:qijin2012@yeah.net

通信地址：410073 湖南省長沙市國防科學技術大學并行與分布處理重點實驗室

Address:National Laboratory of Parallel and Distributed Processing,National University of Defense Technology,Changsha 410073,Hunan,P.R.China