1、投稿日期：2017-03-04.基金项目：E 级计算机关键技术验证系统(2016YFB0200401)；基于张量积的向量化代码自动生成和调优技术研究(61572025)作者简介：张军阳(1987-),男，河南平顶山人，博士研究生， E-mail:郭阳（通信作者），男，研究员，博士，博士生导师，E-mail:面向向量处理器的二维矩阵卷积的设计与实现张军阳，郭阳（国防科技大学 计算机学院，湖南 长沙 410073）摘要：为了加速卷积神经网络模型的计算速度，便于大规模神经网络模型在嵌入式微处理器中的实现，本文以FT-matrix2000 向量处理器体系结构为研究背景，通过对多核向量处理器体系结构的分
2、析和卷积神经网络算法的深入研究，提出一种将规模较小的卷积核数据置于标量存储体，尺寸较大的卷积矩阵置于向量存储体的数据布局方案；针对矩阵卷积中数据难以复用的问题，提出一种根据卷积核移动步长的不同动态可配置的混洗模式，通过对所取卷积矩阵元素进行不同的移位操作，进而大幅提高卷积矩阵数据的复用率；针对二维矩阵卷积由于存在数据相关性进而难以多核并行的问题，提出一种将卷积矩阵多核共享，卷积核矩阵多核独享的多核并行方案；设计了卷积核尺寸不变、卷积矩阵规模变化和卷积矩阵尺寸不变、卷积核规模变化的两种计算方式，并在主流 CPU、gtx1080TI、TI6678、FT-matrix2000 平台进行了性能对比与分
3、析。最终实验结果表明，相比多核 CPU 最高可加速 238 倍，相比 TI6678 可加速21 倍、相比高性能 gtx1080TI 可加速 663805 倍。 关键字：卷积神经网络；向量处理器；多核实现；矩阵卷积 中国分类号： 文献标志码： 文章编号：Design and implementation of two-dimensional matrix convolution based on vector processorZHANG Junyang， GUO Yang(College of Computer, National University of Defense Technolog
4、y, Changsha 410073, China)Abstract: In order to accelerate the computational speed of convolutional neural network model and facilitate the implementation of large-scale neural network model in embedded microprocessor, this paper takes the FT-matrix2000 vector processor architecture as the research
5、background, through the analysis of the multi-core vector processor architecture and convolutional neural network algorithm, a data layout scheme is proposed in which a smaller convolution kernel data is placed in a scalar memory bank and a larger convolution matrix is placed in a vector bank. Aimin
6、g at the problem that the data in the matrix convolution is hard to reuse, a dynamic shuffling pattern with different dynamic configurable parameters based on the moving steps of the convolution kernel is proposed, by carrying out different shift operations on the convolution matrix elements, the mu
7、ltiplexing rate of convolution matrix data is greatly improved. Aiming at the problem that two-dimensional matrix convolution is difficult to multi core parallelism due to the existence of data correlation, a multi-core parallel scheme with convolution matrix sharing and convolution kernel matrix mu
8、lti core exclusive is proposed. Two computing methods of convolution kernel size unchanged, convolution matrix size change and convolution matrix size unchanged and convolution kernel size change are designed, performance comparison and analysis are carried out in mainstream CPU, gtx1080TI, TI6678 a
9、nd FT-matrix2000. The final experimental results show that compared to multi-core CPU can be accelerated up to 238 X, compared to TI6678 can be accelerated 21 X, compared to high-performance gtx1080TI can accelerate 663805 X.Keywords: convolution neural network; vector processor; multi-core implemen
10、tation; matrix convolution近年来，深度学习 1受到了越来越多的关注，尤其是在图像识别 2、语音识别 3、机器翻译 4自然语言处理等领域基于深度学习的目标识别技术都取得了重大突破，进而引发了机器学习和计算机视觉等领域的新一轮研究热潮。深度学习包括一系列神经网络模型，如卷积神经网络(Convolutional Neural Network，CNN) 5-6、深度置信网络(Deep Belief Network，DBN) 7-9、自动编码机 10、循环神经网络(Recurrent Neural Network，RNN) 11等常用模型。尤其是基于CNN 的模型在图像识别领域
11、取得了重大突破，当前几乎所有的图像识别类任务中识别率最好的模型都是基于 CNN。一般来说 CNN 模型有若干卷积层和池化层交替出现，最后由若干全连接层和分类层组成。其中卷积层的计算量占整个模型的 85%以上 12，因此，当前的许多研究大多是针对 CNN 的加速器，如图形处理单元(Graphic Processing Unit，GPU) 13-14、现场可编程逻辑门阵列(Field Programmable Gated Array，FPGA) 15-16、专用集成电路(Application Specific Integrated Circuit，ASIC) 17-19、向量数字信号处理器(Di
12、gital Signal Processor，DSP)等。向量处理器体系结构就是其中的一种新颖体系结构 20。一般包括标量处理单元和向量处理单元，标量计算单元负责标量任务的计算和流控；向量计算单元主要负责大规模的向量计算，它包括若干向量处理单元，每个处理单元上包含丰富的运算部件，具有强大的计算能力，可以大幅提高系统的计算性能，与此同时也对国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY2软件开发提出了新的挑战，如何针对向量处理器多处理单元、多功能部件的体系结构特点，充分开发各个层次的并行性，将各类应用高效地向
13、量化是当前面临的主要困难 21。本文针对卷积神经网络模型中二维矩阵卷积难以并行的难题，结合向量处理器的体系结构特点，设计了一种通过混洗操作来提高算法的并行性和提高数据复用率的方法，针对多输出卷积结果矩阵的计算，提出一种加速矩阵卷积计算的多核实现方案，并从体系结构、多级存储和计算模式等方面进行了算法优化和性能分析。指令派发标量 指令标量处理单元 （ S P U ）S P E全局缓存向量 指令向量处理单元 （ V P U ）V P E 间的混洗网络V P E 0 V P E 1 V P E 1 4 V P E 1 5 向量 L o a d / s t o r e向量数据访问地址产生单元B a n
14、k 0 B a n k 1 5 F M A C 1F M A C 2F M A C 3B PD M A共享寄存器D D R图 1FT-matrix2000处理器体系结构Fig.1 Architecture of FT-matrix20001 向量处理器 FT-matrix2000 的体系结构FT-matrix2000 是由国防科技大学微电子所自主研发的一款面向高密度计算的高性能浮点多核向量处理器，单芯片集成 12 颗向量处理器内核，主频 1GHz，双精度峰值性能达到1152GFLOPS。其单核结构如图 1 所示，每个单核是独立的超长指令字(Very Long Instruction Word，
15、 VLIW)体系结构，包括向量处理部件(Vector Processing Unit，VPU)和标量处理部件(Scalar Processing Unit，SPU) ，SPU 负责标量计算和流控，VPU 负责向量计算，包括 16 个向量处理单元(Vector Processing Element，VPE)，每个 VPE 包含一个局部寄存器文件和 3 个浮点乘累加单元(Floating point Multiply Accumulator， FMAC)、2 个 Load/Store 和 1 个BP 共 6 个并行功能部件。所有 VPE 的同一编号的局部寄存器在逻辑上又组成一个 1024 位的向量
16、寄存器。SPU 和 VPU 之间通过共享寄存器交换数据，支持广播指令将标量寄存器数据广播到向量寄存器。向量处理器核支持 11 发射且长度可变的 VLIW 指令，包括 5 条标量指令和6 条向量指令，指令派发单元对执行包进行识别并派发到相应的功能单元去执行。FT-matrix2000 提供 96KB 的标量存储器(Scalar Memory，SM)和 768KB 的阵列存储器(Array Memory，AM) 用于向量访问；多核采用共享DDR 的存储结构，DDR 与 L1D 间支持多核共享的全局缓存(Global Cache，GC)，方便实现核间数据交换和共享。2 矩阵卷积算法介绍与分析2.1
17、卷积神经网络的基本结构一个具有代表性的卷积神经网络结构如图2 所示，该模型主要用于早期的手写数字的识别任务中，而当前的神经网络模型与该模型大同小异，主要包括卷积层、下采样层、全连接层以及相应的正则化层，只是在模型的深度和广度以及连接方式上进行了改动。由于卷积神经网络模型中的卷积层的计算时间约占整个模型计算量的 85%以上，因此，本文主要研究加速卷积神经网络模型中卷积层的计算。 28 4128844 输入层 卷积层 采样层 卷积层 采样层 输出层图 2典型卷积神经网络结构Fig.2 Typical structure of CNN2.2 二维矩阵卷积二维矩阵卷积常用于图像处理中，卷积神经网络中的
18、卷积操作可以抽象成一组卷积核与一个输入特征图的卷积运算，每一个尺寸为国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY3kxky的卷积核在一个输入特征图上与一个kxky的卷积窗口进行对应点积操作并累加求和后得出输出特征图中的一个元素，卷积核在输入特征图上从上到下、从左到右进行如上操作，最终会得出一个二维的输出特征图，设水平移动步长为 Sx，垂直移动步长为 Sy，那么输出特征图上位置为(a，b)元素的计算如式 (1)：(1)1, , ,0yxxymoKoimioiab jaSbjiAijOf I 其中 是输入特征
19、图#mi 和输出特征图#mo,ij之间的卷积核， 是输入特征图与输出特,io征图之间对应的偏置值， 表示输入特征图m与对应输出特征图连接的集合， 表示非线()f性激活函数，如 tanh、sigmoid 或 ReLU 等。2.3 二维矩阵卷积的常用计算方法1 12 21 11 10 11 01 4 2 01 5 2 41 2 01 1 30 2 20 2 10 3 21 1 01 2 10 1 33 3 21 00 12 12 11 22 01 2 1 1 0 2 0 3 1 2 0 12 0 1 3 2 1 3 2 2 1 1 31 1 0 2 0 3 1 1 0 1 3 31 3 2 2 3
20、 2 1 0 1 3 3 21 11 02 02 11 2 1 11 21 10 11 2 1 20 0=1 4 1 2 2 0 2 41 5 1 7 2 4 2 6输入特征图输入特征图矩阵卷积核矩阵输出特征图矩阵卷积核输出特征图传统卷积计算矩阵乘形式的卷积（ a ）（ b ）1 2 2 41 7 2 6图 3 二维矩阵卷积计算的两种方式Fig.3 Two methods of two-dimensional matrix convolution for calculating图 3(a)为传统的矩阵卷积计算方式，输入为 3 个 33 的输入特征图，输入卷积核为 2 个3 通道的 22 卷积核
21、，通过卷积计算可以得出2 个 22 的输出特征图，我们可以发现如果采用传统矩阵卷积计算方式，由于都是二维的矩阵卷积计算，尤其是当卷积核规模较小的时候，算法的并行性难以得到保证。因此，北京邮电大学 ZHANG Qiang-Qiang 等人 22使用了一种新的计算方法，即将卷积矩阵和卷积核矩阵展开成矩阵块的方式进而通过普通矩阵相乘的方式进行卷积计算，如图 3(b)所示。其计算步骤如下：根据卷积核规模和移动步长，将对应的卷积矩阵相应的卷积块展开成列向量的形式，所有展开的列向量组成一个新的矩阵；将单个卷积核矩阵按行展开成一个行向量，所有展开的卷积核组成一个新的卷积核矩阵；将新的卷积矩阵和新的卷积核矩阵
22、做普通矩阵与矩阵乘法，进而可以通过调用普通的 BLAS(Basic Linear Algebra Subprograms，BLAS) 23函数库中的GEMM(General Matrix Multiplication，GEMM)完成二维矩阵卷积的计算。通过图 3(b)所示的方法将小的难以并行的卷积计算模块组合成大规模矩阵的形式，进而利用矩阵乘法的高并行性来提高矩阵卷积的计算效率。由于本文研究的矩阵卷积操作主要应用于卷积神经网络模型中，且卷积计算的结果通常要作为下一级运算的输入值，因此，神经网络模型中卷积层的计算结果有特定的二维结构，而图 3(b)中采用的计算方式，虽然通过组合成大规模矩阵的形式
23、增加了矩阵卷积计算的并行性，但是同时也破坏了卷积计算结果的二维结构，从而使得该计算结果不能直接作为下一级的输入值，可能需要对该卷积结果进行数据重排序才能作为下一级操作的输入值。同时，由于是对卷积矩阵进行了有重复展开，因此展开后的矩阵规模大大增加了，如 33 的卷积矩阵在 22 的卷积核规模下的展开为 44，数据量增加了 77.8%，而当卷积矩阵为 55 时，展开的卷积矩阵变成了 165，数据量增加到了 220%，而片上存储对计算来说尤为重要，因此该方法使得数据存储量的增加是该方法的一大缺陷。3 二维矩阵卷积的向量化设计与实现通过对第 2 节传统卷积计算方法的分析，我们提出一种新的二维矩阵卷积的
24、向量化实现方法，该算法实现的流程如下：1）根据多核向量处理器的核数 m，输入特征图的数量 n，尺寸为 n1n2，输入卷积核的数量 k，尺寸为 k1k2，向量处理器单核向量存储体的容量 v1和标量存储体的容量 v2，合理划分多核处理器中每个核的计算负载；2）将输入特征图元素按行连续存储于向量存储体 AM 中，卷积核矩阵按行连续存储于标量存储体 SM 中；3）根据单核向量处理单元的个数加载 p 个输入特征图元素，标量加载卷积核第 1 个元素，并广播至向量寄存器中；4）将步骤 3）中通过广播得到的向量寄存器与向量加载的 p 的输入特征图元素进行乘加操作，并将计算结果累加至累加寄存器 ACC 中；5）
25、根据卷积核在输入特征图中的移动步长，配置相应的混洗模式，并将步骤 3）中所加载的p 个输入特征图元素使用该混洗模式进行移位操作，同时标量顺序加载第二个卷积核元素，并广播至向量寄存器中，将该寄存器与移位后的寄存器中的元素进行乘法操作，并将乘法结果累加至步骤 4）中的累加寄存器 ACC 中；6）根据卷积核的行数 k1和列数 k2重复上述过程，直到完成输出特征图第 1 行元素的计算；国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY47）根据输入特征图的行数 n1和列数 n2，向量加载输入特征图的下一行元素，重复上述
26、步骤，进而完成整个卷积结果矩阵的计算。卷积神经网络模型中的卷积核规模一般都比较小，通常为11、33、55、77、1111 等，根据我们的算法需要，卷积核数据需要标量取，因此将卷积核数据置于 SM，输入特征图矩阵置于AM。其向量化实现过程如下：3.1 单核程序向量化方法1)我们以双精度浮点为例，单个数据为 8 字节，为了叙述方面我们设输入特征图为 NN的方阵，则 NN8=7681024，所以向量存储体 AM 中单次可以存放的最大输入特征图尺寸为 222222，对于嵌入式处理器编程来说，当输入特征图的尺寸大于 222222 时，我们就需要从外存 DDR 中通过 DMA 来加载数据。1,kb,12,
27、3kb1,4,15kb,6#广 播A C C1A C C2A C C3A C C4A C C5A C C6累 加 ,2,1,21,1,20广 播A C C1A C C2A C C3A C C4A C C5A C C6累 加12,kb,1kb2,12kb,#3广 播A C C1A C C2A C C3A C C4A C C5A C C6累 加 2,2,2,04广 播A C C1A C C2A C C3A C C4A C C5A C C6累 加 = = = = =O U T1O U T2O U T3O U T4O U T5图 4 FT-matrix2000 矩阵卷积的算法映射( B：66；步长：1
28、1；kernel：22)Fig.4 Algorithm mapping of FT-matrix2000(B:66;stride:11;kernel:22)我们以一个例子来说明在不展开卷积矩阵的情况下如何同时计算输出卷积矩阵多个元素的向量化实现方法，考虑一个 16 的PEs(PE1，PE 2， PE3，PE 4，PE 5，PE 6)，输入特征图矩阵 I 为 66，卷积核矩阵为 22，步长为 1，为了叙述简洁，我们只描述输出结果矩阵第 1 行元素的计算(如图 4 所示，其他行计算过程类似) 。#1：6 个 PEs 加载输入矩阵 B的第 1 行元素(b1,1，b 1,2，b 1,3，b 1,4，b
29、 1,5，b 1,6)，标量加载卷积核矩阵的第 1 行第 1 个元素(k 1,1)，将 k1,1广播至向量寄存器中，即向量寄存器中的元素为k1,1，k 1,1，k 1,1，k 1,1，k 1,1，k 1,1，通过向量处理器的乘加指令 VFMULAD，每一个 PE 完成对应元素与 k1,1的乘法，并将乘法结果累加至对应的累加寄存器 ACC 中；#2：根据卷积核的移动步长配置相应的混洗模式( 如图 5 所示)，本实例中由于移动步长为1，因此配置的混洗模式为对应的向量寄存器中的元素向左移动 1 位，最左边的元素移出寄存器，最右边的元素补 0，移位后当前 PEs 所对应的元素为 b1,2，b 1,3，
30、b 1,4，b 1,5，b 1,6，0，同时加载卷积核的第 1 行第 2 个素 k1,2并广播至向量寄存器中，即为 k1,2，k 1,2，k 1,2，k 1,2，k 1,2，通过向量乘加指令，每一个 PE 完成对应元素与 k1,2的乘法，并将乘法结果累加至#1 中的累加寄存器 ACC 中；#3：6 个 PEs 加载输入矩阵 B的第 2 行元素(b2,1，b 2,2，b 2,3，b 2,4，b 2,5，b 2,6)，标量加载卷积核矩阵第 2 行第 1 个元素(k 2,1)并广播至向量寄存器中，通过向量乘加指令，所有 PE 完成对应元素与卷积核元素 k2,1的乘法操作，并将乘法结果累加至#2 中的
31、累加寄存器中；#4：同#2 中的混洗模式，将#3 中取到的元素进行移位，完成移位后当前 PEs 所对应的元素为 b2,2，b 2,3，b 2,4，b 2,5，b 2,6，0，同时标量加载卷积核矩阵第 2 行第 2 个元素并广播至向量寄存器中，即为k2,2，k 2,2，k 2,2，k 2,2，k 2,2，k 2,2，使用向量乘加指令，所有 PE 完成对应元素与 k2,2的乘法操作，并将结果累加至#3 中的累加寄存器中 ACC。完成以上步骤，则#4 中累加寄存器中的值即为卷积结果矩阵第 1 行元素。B1 , 1B1 , 2B1 , 3B1 , 4B1 , 5B1 , 6B1 , 2B1 , 3B1
32、 , 4B1 , 5B1 , 60左 移图 5 步长为 1 的混洗操作Fig.5 Shuffle mode of stride 1通过对本文提出的矩阵卷积的向量化实现方法分析，可以看出本方法有以下优势，不需要通过将输入特征图和卷积核进行重组成大规模的矩阵来提高并行性，减轻了矩阵展开所带来了存储压力，提高了所取数据的利用率，在向量化代码的优化过程中，我们可以通过软件流水的方式提高算法的并行性，完全流水的情况下，每拍可以同时计算出输出特征图的一行元素，同时也避免了文献15中所采用的 PEs 阵列间数据的频繁移动操作。传 输A r r a y M e m o r yA M0A M1D M AD M
33、AD D RD D R传 输A r r a y M e m o r yA M0A M1D M AD M AD D RD D R图 6 双 Buffer 机制示意国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY5Fig.6 Double buffer mechanism2)当输入特征图尺寸超过向量存储体的容量时，此时，向量存储体不能一次加载整个输入特征图。由于在高性能计算中，需要平滑各级存储之间的数据搬移，保证内核计算所需要的数据供给，进而取得高效的处理器效率。由于FT-matrix2 拥有多级存储结构，如寄存
34、器文件、L1D、Array Memory、Global Cache 和 DDR，且 L1D 可以配置为全 Cache 和全 SRAM 两种模式，因此，当输入特征图数量超过向量存储体的容量时，我们采用 SRAM 的方式，将 L1D 配置为全 SRAM 的形式，采用双缓冲的乒乓方式平滑多级存储结构的数据传输，使得内核的计算与 DMA 数据搬移完全重叠，让内核始终处于峰值运行状态，进而取得最佳计算效率(双缓冲机制如图 6 所示)。值得注意的是由于卷积计算有数据的重叠，因此在卷积矩阵分块传输时，根据卷积核的尺寸，第 2 次传输的卷积矩阵需要有部分数据与第一次传输的数据重叠，假设卷积核尺寸为mm，若第
35、1 次传输 1 到 n 行，则第 2 次需重叠传输第 1 次传输的后 m-1 行元素，由于卷积核一般较小，因此重叠传输数据的比重并不大。3.2 矩阵卷积的多核程序设计分析算法的并行不仅有单核内部的向量化实现，也包括多核的核间并行，由于 FT-matrix2000 是一款多核处理器，因此，研究矩阵卷积的多核并行也是一个重要内容，尤其是对于大规模的矩阵卷积来说。由于核间并行需要考虑计算负载的多核划分、算法的相关性、核间的通信开销等，因此多核程序的设计更加复杂。由于当前卷积神经网络模型中卷积层的计算规模从11 到 512512 都存在，且一般都是卷积核尺寸较小的矩阵卷积操作，因此，在矩阵卷积的多核实
36、现中一般不会将单个输入特征图分配给多个核来计算，而是由每个核负责独立的卷积计算，即通过将输入特征图多核共享，卷积核矩阵多个核独有的方式来进行矩阵卷积的多核实现。核 2卷积核缓存输入特征图卷积核核 1核 n输入缓存图 7 FT-matrix2000 矩阵卷积多核实现方案Fig.7 Multicore implementation of matrix convolution on FT-matrix2000此外，由于在矩阵卷积的计算过程中，因卷积核的尺寸以及卷积核在输入特征图上的滑动步长不同，单个卷积的计算过程存在一定的数据相关性，因此，考虑到多核同步的通信开销，单输入特征图的卷积计算不易进行多核
37、并行。因此本文主要针对单核实现单个卷积矩阵的向量化设计，当有多个独立的卷积计算时再由多核来并行进行加速。多输出矩阵卷积的多核实现如图 7 所示，将 DDR 中的输入特征图广播至 n 个计算核心，同时 n 个核加载不同的卷积核至 n 个核的标量存储体 SM 中，则多核程序实现中每个核的计算过程同单核程序一样。4 性能测试与分析本实验的对比平台有 AMD A10-6700 APU，4 核 4 线程， DDR3 16G 内存，主频3.7GHz，GPU 平台为 NVIDIA GeForce GTX1080TI，11GB 显存，核心频率 1.5GHz，使用 NVIDIA 专门为卷积计算高度优化的cudn
38、n24函数库，并基于 Torch25神经网络编程框架进行矩阵卷积的性能统计和优化，FT-matrix2000 平台单核核内向量存储空间768KB，标量存储空间 96KB，核外 DDR 最大支持 128GB，主频 1GHz；使用 TI 的高性能多核数字信号处理器 TMS320C6678，8 核，1.25GHz 主频，每个核拥有 32KB 的 L1P 和32KB 的 L1D，并基于 CCS5.5 软件编程平台完成所有程序代码的测试与性能统计，飞腾平台使用类 CCS 的 FT-matrix2000 软件编程平台并完成 FT-matrix2000 向量程序代码的编写与性能统计，核心代码程序采用基于手工
39、汇编的方式进行软件流水和循环展开优化。通过 2.2 节中的分析可知二维矩阵卷积是典型的计算密集型和访存密集型操作，而当前的卷积神经网络模型中矩阵卷积的计算又占据着整个模型计算量的 85%以上，因此本实验选取两种计算模式进行分析。第 1 种，卷积矩阵不变，卷积核矩阵规模相应增加；第 2 种，卷积核矩阵不变，卷积矩阵规模相应增加。(图中所有类似 165 都表示 1616，55 的方阵，且都为 64 位双精度浮点)。国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY6图 8 基于卷积核规模变化的 FT-matrix2
40、000 与 CPU/6678的加速比Fig.8 Speedup of FT-matrix2000 and CPU/6678 based on kernel scale change图 8 统计了随着卷积核尺寸变化，FT-matrix2000 对于 CPU 和 TI6678 的加速比，可以看出 FT-matrix2000 向量处理器相对于主流多核CPU 取得了 2444 倍不等的加速比，相对于TI6678 取得了 1417 倍不等的加速比。图 9 为卷积核规模不变，随着输入矩阵规模变化的 FT-matrix2000 基于 CPU 和 TI6678 的加速比，在此计算模式下相对主流 CPU 取得了
41、 14238 倍不等的加速比，相对 TI6678 取得了 1321 倍不等的加速比。图 9 基于卷积矩阵变化的 FT-matrix2000 与 CPU/6678的加速比Fig.9 Speedup of FT-matrix2000 and CPU/6678 based on matrix scale change图 8 和图 9 分别为两种卷积计算模式下，FT-matrix2000 向量处理器相对于主流 CPU 和TI6678 取得的性能优势，由于我们的算法主要是用于卷积神经网中卷积层的计算，因此，采取的数据计算规模不大，从图中可以看出当矩阵卷积计算的规模较小时可以取得较大的加速比，而当矩阵规模
42、增加时加速比有所下降，并基本保持稳定。可见，此两种计算模式，对程序的实现性能有一定的影响。图 10 基于卷积核规模变化的 FT-matrix2000 相对 GPU的加速比Fig.10 Speedup of FT-matrix2000 and GPU based on kernel scale change 图 11 基于输入特征图规模变化的 FT-matrix2000 相对GPU 的加速比Fig.11 Speedup of FT-matrix2000 and GPU based on matrix scale change 图 10 为 FT-matrix2000 在基于卷积核规模变化的模式下
43、基于 GPU 取得了 34882 至 119523倍不等的加速比。图 11 为 FT-matrix2000 在基于输入特征图规模变化的模式下相对 GPU 取得了 3198 至 663805 倍不等的加速比。从图 10 中可以看出，当输入特征图矩阵不变时，随着卷积核矩阵规模的增加，FT-matrix2000 相对 GPU的加速比呈下降趋势，主要因为 FT-matrix2000属于嵌入式向量处理器，而 NVIDIA 1080TI 属于服务器级高性能处理器，拥有 3584 个 cuda核心，单精度峰值性能高达 11.5TFLOPS，但当矩阵卷积规模较小的时候，GPU 由于计算负载不足，其并没有完全发
44、挥出应有的计算性能，同时由于在使用 GPU 计算的时候，需要通过使用 cudaMemcpy()函数将需要计算的数据由 CPU内存拷贝到 GPU 显存，因此，当数据量较小的时候数据的拷贝要占据大部分时间，影响了计算效率，故加速比很高，而当计算规模增加的时候，加速比下降，主要因为数据拷贝占用时间比重降低，计算占比增加，此时数据量越大越能发挥 GPU 的众核优势。图 11 中的整体趋势同图 10 类似，只是加速比并没有图 10 高，主要是因为 FT 平台的算法实现中，内核循环是由卷积核的尺寸控制，当卷积核尺寸越大，内核程序通过软件流水所获得的性能优势越明显，因此，FT 平台该模式下矩阵卷积的算法实现
45、效国 防 科 技 大 学 学 报JOURNAL OF NATIONAL UNIVERSITY OF DEFENSE TECHNOLOGY7率越高。针对卷积核尺寸对算法实现效率影响较大的特点，下一步可考虑能否在不增加卷积矩阵规模的前提下，通过融合多个卷积核矩阵来提高 FT-matrix2000 向量处理器计算矩阵卷积的计算效率。5 结论针对卷积神经网络中二维矩阵卷积计算中存在的卷积矩阵和卷积核规模较小、难以向量化和处理器计算资源利用率不高的问题。本文提出一种基于混洗的二维矩阵卷积向量化实现方法，通过配置不同的混洗模式可以实现卷积核不同移动步长的卷积计算，同时通过对输入特征图数据的移位操作可以充分
46、利用已取数据，大大提高数据的利用率，针对大规模的矩阵卷积提出了一种多核实现方案，采用 L1D 配置的SRAM 模式，用双缓冲的乒乓方式平滑多级存储结构的数据传输，使得内核的计算与 DMA 数据搬移完全重叠，进而提高算法的计算效率。最后，基于主流 CPU、TI6678、高性能 GPU 进行了算法的性能统计和分析，实验结果显示了FT-matrix2000 比 CPU、TI6678 及 GPU 具有更好的计算优势，相比 CPU 最高可加速 238 倍，相比 TI6678 可加速 21 倍，相比 GPU 可加速663805 倍。本文研究表明，二维矩阵卷积可在FT-matrix2000 向量处理器上可以
47、取得良好的性能，为今后国产 FT-matrix2000 高性能多核向量处理器面向边缘计算 26做了很好的准备。参考文献(Reference)1 Deng L, Yu D. Deep Learning: Methods and ApplicationsM. Now Publishers Inc. 2014.2 Simonyan K, Zisserman A. Very Deep Convolutional Networks for Large-Scale Image RecognitionJ. Computer Science, 2014.3 Graves A, Jaitly N. Towards end-to-end speech recognition with recu
