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