一种基于嵌入式GPU和卷积计算的图像处理方法和装置与流程

本发明涉及计算机视觉技术领域，更具体的说是涉及一种基于嵌入式gpu和卷积计算的图像处理方法和装置。

背景技术：

自从在imagenet竞赛中引入卷积神经网络以来，计算机视觉技术在过去几年中取得了长足的进步，在诸如图像分类、模式识别和多媒体压缩等各种领域表现出了很好的性能。其中ssd算法得到了广泛的应用，ssd算法均匀地在图片的不同位置进行密集抽样，抽样时可以采用不同尺度和长宽比，然后利用卷积神经网络提取特征后直接进行分类与回归，整个过程只需要一步，速度快于rcnn系列算法。ssd算法在多尺度特征图、利用卷积进行检测、设置检验框等方面做了优化，更适合检测相对较小的目标。

由于ssd检测算法过程比较复杂，当在嵌入式硬件平台上实现时，需要消耗大量的存储、计算单元，这就对硬件平台提出了较高的要求。由于用dsp和fpga等硬件编程相对于软件开发来说难度很大，另一方面很多针对软件的算法很难以用硬件实现，而且硬件开发的周期更长、成本更高，故综合考虑选用嵌入式gpu作为处理器。嵌入式gpu可以并发执行程序、支持深度学习cuda库；然而，嵌入式gpu内存有限，因此，如何在嵌入是平台上优化卷积计算的内存利用率和运行时间对于图像处理就显得十分重要。

因此，如何提供一种基于嵌入式gpu和卷积计算的图像处理方法是本领域技术人员亟需解决的问题。

技术实现要素：

有鉴于此，本发明提供了一种基于嵌入式gpu和卷积计算的图像处理方法和装置，能够降低内存开销、提高计算效率。

为了实现上述目的，本发明采用如下技术方案：

一种基于嵌入式gpu和卷积计算的图像处理方法，包括：

s1：采用内存优化的卷积扩展方法对输入图像进行矩阵变换和cuda并行处理，得到中间矩阵；

s2：对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；

s3：对中间矩阵和卷积核子矩阵进行卷积计算；

s4：对卷积计算后输出的子矩阵进行合并，得到输出矩阵。

优选的，在步骤s1中，采用内存优化的卷积扩展方法对输入图像进行矩阵变换的具体步骤包括：

s11：在输入图像对应的输入矩阵i[0:ih,0:iw]上，选取[0:ih,0:sw]为滑动窗口，滑动窗口大小为ih*sw；ih为输入矩阵的行，iw为输入矩阵的列；

s12：滑动窗口[0:ih,0:sw]依次在输入矩阵i[0:ih,0:iw]上开始滑动，滑动步长为1，滑动窗口内的数据按行拉伸为中间矩阵l[0:oh,0:ow]的一行；

s13：滑动窗口[0:ih,x:sw]在输入矩阵i[0:ih,0:iw]上滑动iw-sw+1个单位后结束，形成中间矩阵l[0:oh,0:ow]，oh＝iw-sw+1，ow＝ih*sw，中间矩阵的行数为iw-sw+1，oh为中间矩阵的行，ow为中间矩阵的列。

优选的，步骤s2的具体步骤包括：

s21：以卷积核矩阵s[0:sw*sh，1]为滑动窗口，滑动窗口在卷积核临时矩阵s上滑动，其中，卷积核临时矩阵s初始化为空；sw为卷积核矩阵的长，sh为卷积核矩阵的宽；

s22：滑动窗口s在卷积核临时矩阵s第1列上向下滑动0个单位，s内其余元素补0，形成卷积核临时矩阵s的第1列；

s23：滑动窗口s在卷积核临时矩阵s第i列上向下滑动3*(i-1)个单位，s内其余元素补0，形成卷积核临时矩阵s的第i列；

s24：基于上述步骤得到卷积核临时矩阵s[0:sh,0:sw]，且卷积核临时矩阵s[0:sh,0:sw]中每列有效数据与卷积核矩阵s[0:sw*sh，1]的数据相对应，s上其余元素为0；sh＝ih*sw，sw＝ih-sh+1；sh为卷积核临时矩阵的行，sw为卷积核临时矩阵的列；

s25：对卷积核临时矩阵s[0:sh,0:sw]做分块处理，分为卷积核第一子矩阵s1[0:sh,0:sw1]和卷积核第二子矩阵s2[0:sh,0:sw2]，其中sw1＝sw-sw/2，sw2＝sw/2；，sw1为分块后卷积核第一子矩阵的列，sw2为分块后卷积核第二子矩阵的列。

优选的，在步骤s3中，利用cuda库中高度优化的cublas函数对中间矩阵和卷积核子矩阵进行卷积计算。

优选的，步骤s3的具体步骤包括：

s31：创建并初始化cublas库对象；

s32：在gpu中为待运算的数据以及需要存放结果的变量开辟显存空间；

s33：将待运算的数据从内存传输进显存；

s34：基于cublas函数对中间矩阵l[0:oh,0:ow]分别和卷积核第一子矩阵s1[0:sh,0:sw1]、卷积核第二子矩阵s2[0:sh,0:sw2]进行卷积计算；

s35：从gpu中获取卷积计算后的第一子矩阵o1[0:oh,0:ow1]和第二子矩阵o2[0:oh,0:ow2]，释放开辟的显存空间以及cublas库对象。

优选的，步骤s4具体包括：

将第一子矩阵o1[0:oh,0:ow1]和第二子矩阵o2[0:oh,0:ow2]合并为输出矩阵o[0:oh,0:ow]，其中，oh为输出矩阵的行，ow为输出矩阵的列。

优选的，该方法由cpu和gpu协同处理实现，其中，gpu负责执行步骤s1和步骤s3，cpu负责执行步骤s2和步骤s4。

一种基于嵌入式gpu和卷积计算的图像处理装置，包括：

卷积扩展模块，用于采用内存优化的卷积扩展方法对输入图像进行矩阵变换和cuda并行处理，得到中间矩阵；

扩充分块处理模块，用于对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；

卷积计算模块，用于对中间矩阵和卷积核子矩阵进行卷积计算；

合并模块，用于对卷积计算后输出子矩阵进行合并，得到输出矩阵。

经由上述的技术方案可知，与现有技术相比，本发明公开提供了一种基于嵌入式gpu和卷积计算的图像处理方法和装置，针对ssd算法中的卷积计算进行优化，采用内存优化的卷积扩展对输入图像进行矩阵变换，利用cuda并行处理形成中间矩阵，同时采用卷积核矩阵行列扩充对齐，卷积核矩阵扩充后做分块处理以减少运算时内存开销，最后采用cuda库中高度优化的cublas矩阵乘法函数进行卷积计算并行加速，最后合并输出矩阵。

与现有技术相比具有如下优点：

1、本发明采用的方法，与im2col卷积计算方法相比，减少了内存开销，提高了卷积计算速度。在嵌入式gpu平台jetsontx2上，经大量实验证明，本发明提供的方法平均内存使用效率提升了45％，卷积计算速度平均提高了90％以上。

2、本发明采用的卷积核矩阵扩展使其与中间矩阵地址对齐，可充分简化分段卷积控制逻辑，减少cpu与gpu之间的数据传输次数，进而节省gpu总线传输资源。

3、本发明基于嵌入式gpu平台，可充分利用cuda库，发挥gpu并行控制的优势，从而加速cnn卷积计算的运行。适合运行在内存受限的平台，如嵌入式gpu、物联网等，而且适合cnn和dnn的应用。

因此，综上所述，本发明提供的基于嵌入式gpu和卷积计算的图像处理方法和装置具有很好的推广意义。

附图说明

为了更清楚地说明本发明实施例或现有技术中的技术方案，下面将对实施例或现有技术描述中所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本发明的实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据提供的附图获得其他的附图。

图1为本发明提供的基于嵌入式gpu和卷积计算的图像处理方法的流程图；

图2为本发明提供的图像输入像素的卷积扩展示意图；

图3为本发明提供的卷积核扩展及分块示意图；

图4.1为本发明提供的cublas矩阵相乘示意图一；

图4.2为本发明提供的cublas矩阵相乘示意图二；

图5为本发明提供的输出子矩阵合并示意图；

图6为本发明提供的实验测试基准集；

图7为本发明提供的实验结果图。

具体实施方式

下面将结合本发明实施例中的附图，对本发明实施例中的技术方案进行清楚、完整地描述，显然，所描述的实施例仅仅是本发明一部分实施例，而不是全部的实施例。基于本发明中的实施例，本领域普通技术人员在没有做出创造性劳动前提下所获得的所有其他实施例，都属于本发明保护的范围。

参见附图1，本发明实施例公开了一种基于嵌入式gpu和卷积计算的图像处理方法，包括：

s1：采用内存优化的卷积扩展方法对输入图像进行矩阵变换和cuda并行处理，得到中间矩阵；

s2：对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；

s3：对中间矩阵和卷积核子矩阵进行卷积计算；

s4：对卷积计算后输出的子矩阵进行合并，得到输出矩阵。

为了进一步优化上述技术方案，在步骤s1中，采用内存优化的卷积扩展方法对输入图像进行矩阵变换的具体步骤包括：

s11：在输入图像对应的输入矩阵i[0:ih,0:iw]上，选取[0:ih,0:sw]为滑动窗口，滑动窗口大小为ih*sw；ih为输入矩阵的行，iw为输入矩阵的列；

s12：滑动窗口[0:ih,0:sw]依次在输入矩阵i[0:ih,0:iw]上开始滑动，滑动步长为1，滑动窗口内的数据按行拉伸为中间矩阵l[0:oh,0:ow]的一行；

s13：滑动窗口[0:ih,x:sw]在输入矩阵i[0:ih,0:iw]上滑动iw-sw+1个单位后结束，形成中间矩阵l[0:oh,0:ow]，oh＝iw-sw+1，ow＝ih*sw，中间矩阵的行数为iw-sw+1，oh为中间矩阵的行，ow为中间矩阵的列。

为了进一步优化上述技术方案，步骤s2的具体步骤包括：

s21：以卷积核矩阵s[0:sw*sh，1]为滑动窗口，滑动窗口在卷积核临时矩阵s上滑动，其中，卷积核临时矩阵s初始化为空；sw为卷积核矩阵的长，sh为卷积核矩阵的宽；

s22：滑动窗口s在卷积核临时矩阵s第1列上向下滑动0个单位，s内其余元素补0，形成卷积核临时矩阵s的第1列；

s23：滑动窗口s在卷积核临时矩阵s第i列上向下滑动3*(i-1)个单位，s内其余元素补0，形成卷积核临时矩阵s的第i列；

s24：基于上述步骤得到卷积核临时矩阵s[0:sh,0:sw]，且卷积核临时矩阵s[0:sh,0:sw]中每列有效数据与卷积核矩阵s[0:sw*sh，1]的数据相对应，s上其余元素补0；sh＝ih*sw，sw＝ih-sh+1；sh为卷积核临时矩阵的行，sw为卷积核临时矩阵的列；

s25：对卷积核临时矩阵s[0:sh,0:sw]做分块处理，分为卷积核第一子矩阵s1[0:sh,0:sw1]和卷积核第二子矩阵s2[0:sh,0:sw2]，其中sw1＝sw-sw/2，sw2＝sw/2；，sw1为分块后卷积核第一子矩阵的列，sw2为分块后卷积核第二子矩阵的列。

为了进一步优化上述技术方案，在步骤s3中，利用cuda库中高度优化的cublas函数对中间矩阵分别和卷积核子矩阵进行卷积计算。

为了进一步优化上述技术方案，步骤s3的具体步骤包括：

s31：创建并初始化cublas库对象；

s32：在gpu中为待运算的数据以及需要存放结果的变量开辟显存空间；

s33：将待运算的数据从内存传输进显存；

s34：基于cublas函数对中间矩阵l[0:oh,0:ow]分别和卷积核第一子矩阵s1[0:sh,0:sw1]、卷积核第二子矩阵s2[0:sh,0:sw2]进行卷积计算；

s35：从gpu中获取卷积计算后的第一子矩阵o1[0:oh,0:ow1]和第二子矩阵o2[0:oh,0:ow2]，释放开辟的显存空间以及cublas库对象。

为了进一步优化上述技术方案，步骤s4具体包括：

将第一子矩阵o1[0:oh,0:ow1]和第二子矩阵o2[0:oh,0:ow2]合并为输出矩阵o[0:oh,0:ow]，其中，oh为输出矩阵的行，ow为输出矩阵的列。

为了进一步优化上述技术方案，该方法由cpu和gpu协同处理实现，其中，gpu负责执行步骤s1和步骤s3，cpu负责执行步骤s2和步骤s4。

此外，本发明实施例还公开了一种基于嵌入式gpu和卷积计算的图像处理装置，包括：

卷积扩展模块，用于采用内存优化的卷积扩展方法对输入图像进行矩阵变换和cuda并行处理，得到中间矩阵；

扩充分块处理模块，用于对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；

卷积计算模块，用于对中间矩阵和卷积核子矩阵进行卷积计算；

合并模块，用于对卷积计算后输出子矩阵进行合并，得到输出矩阵。

下面结合具体实例对本发明所提供的技术方案做进一步阐述。

1、将输入图像转换成中间矩阵

(1)输入矩阵i如图2所示，在输入矩阵i[0:7,0:7]上选取i[0:7,0:3]为滑动窗口，滑动窗口的大小为3×7，滑动窗口内的数据按行拉伸为中间矩阵l[0:5,0:21]的第1行，即l[0,0:21]＝{0,1,0,0,2,1,0,0,1,1,0,0,1,1,0,2,1,0,0,0,0}。

(2)滑动窗口以步长为1向前滑动，形成为滑动窗口[0:7,1:4]，滑动窗口内的数据按行拉伸为中间矩阵l[0:5,0:21]的第2行，即l[1,0:21]＝{1,0,1,2,1,0,0,1,2,0,0,1,1,0,0,1,0,0,0,0,0}。

(3)滑动窗口以步长为1向前滑动，形成为滑动窗口[0:7,2:5]，滑动窗口内的数据按行拉伸为中间矩阵l[0:5,0:21]的第3行，即l[2,0:21]＝{0,1,0,1,0,0,1,2,1,0,1,2,0,0,1,0,0,2,0,0,0}。

(4)滑动窗口以步长为1向前滑动，形成为滑动窗口[0:7,3:6]，滑动窗口内的数据按行拉伸为中间矩阵l[0:5,0:21]的第4行，即l[3,0:21]＝{1,0,2,0,0,1,2,1,0,1,2,0,0,1,1,0,2,1,0,0,0}。

(5)滑动窗口以步长为1向前滑动，形成为滑动窗口[0:7,4:7]，滑动窗口内的数据按行拉伸为中间矩阵l[0:5,0:21]的第5行，即l[4,0:21]＝{0,2,1,0,1,0,1,0,0,2,0,1,1,1,0,2,1,0,0,0,0}。

(6)形成中间矩阵l[0:5,0:21]，如图2所示。

2、对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵；

为使其与中间矩阵l[0:5,0:21]地址对齐，将卷积核矩阵s[9,1]扩充为卷积核临时矩阵s[0:21,0:5]，如图3所示。以卷积核矩阵s为滑动窗口；滑动窗口开始在卷积核临时矩阵s上滑动，s上窗口内的数据为卷积核矩阵s上的数据，即卷积核临时矩阵s[0:21,0:5]中每列有效数据与卷积核矩阵s[9,1]＝{1,2,0,1,1,0,1,-1,0}的数据相对应，s上其余元素补0。其具体步骤如下：

(1)滑动窗口开始在s第1列上向下滑动0个单位，形成卷积核临时矩阵s的第1列：s[0:21,0]＝{1,2,0,1,1,0,1,-1,0,0,0,0,0,0,0,0,0,0,0,0,0}。

(2)滑动窗口在s第2列上向下滑动3个单位，形成卷积核临时矩阵s的第2列：s[0:21,1]＝{0,0,0,1,2,0,1,1,0,1,-1,0,0,0,0,0,0,0,0,0,0}。

(3)滑动窗口在s第3列上向下滑动6个单位，形成卷积核临时矩阵s的第3列：s[0:21,2]＝{0,0,0,0,0,0,1,2,0,1,1,0,1,-1,0,0,0,0,0,0,0}。

(4)滑动窗口在s第4列上向下滑动9个单位，形成卷积核临时矩阵s的第4列：s[0:21,3]＝{0,0,0,0,0,0,0,0,0,1,2,0,1,1,0,1,-1,0,0,0,0}。

(5)滑动窗口在s第5列上向下滑动12个单位，形成卷积核临时矩阵s的第5列：s[0:21,4]＝{0,0,0,0,0,0,0,0,0,0,0,0,1,2,0,1,1,0,1,-1,0}。

(6)形成卷积核临时矩阵s[0:21,0:5]，如图3所示。

(7)卷积核临时矩阵s[0:21,0:5]做分块处理，分为子矩阵s1[0:21,0:3]和子矩阵s2[0:21,0:2]，如图3所示。

3、调用cublas函数进行卷积计算

cublas库是cuda专门用来解决线性代数运算的库，支持矩阵乘法等运算。使用cublas库函数，可充分发挥gpu并行执行的优势，加速矩阵的运算。如图4.1、图4.2所示，调用cublas函数进行卷积计算的具体步骤如下：

(1)创建并初始化cublas库对象；

(2)在gpu中为待运算的中间矩阵l[0:5,0:21]、卷积核第一子矩阵s1[0:21,0:3]和卷积核第二子矩阵s2[0:21,0:2]以及需要存放结果的输出子矩阵o1[0:5,0:3]和o2[0:5,0:2]开辟显存空间；

(3)调用cudamemcpy()等函数将待运算的中间矩阵l[0:5,0:21]以及卷积核子矩阵s1[0:21,0:3]、卷积核子矩阵s2[0:21,0:2]传输进显存；

(4)调用cublas库函数cublassgemm()进行对中间矩阵l[0:5,0:21]和卷积核子矩阵s1[0:21,0:3]和卷积核子矩阵s2[0:21,0:2]分别进行相乘；

(5)从gpu中获取结果变量输出子矩阵o1[0:5,0:3]和o2[0:5,0:2]；

(6)释放申请的显存空间以及cublas库对象。

4、把第一子矩阵o1[0:5,0:3]和第二子矩阵o2[0:5,0:2]合并为输出矩阵[0:5,0:5]，如图5所示。

下面结合实验结果对本发明提供的技术方案做进一步说明。

为了进行比较，对im2col卷积计算的优化前后结果做比较，本发明建立了一个全面的基准集，如图6所示。该基准集包括11个独特的卷积层，来自各种公共的卷积神经网络。实验在嵌入式gpu平台nvidiajetsontx2上进行，以gpu的时钟频率为基准，对程序的运行时间和内存进行了测量，每个算法运行10次，并取平均值，实验结果如7所示。

通过实验结果可知，优化后的卷积计算在运行时间和运行内存上，与优化前相比性能有了大幅度提升，内存开销大大减少，运行时间得到提高。特别对于ssd算法的卷积层，3×3卷积，图像输入分辨率在300×300的情况下，运算时间提高了90％以上，运行内存减小了45％以上。

本说明书中各个实施例采用递进的方式描述，每个实施例重点说明的都是与其他实施例的不同之处，各个实施例之间相同相似部分互相参见即可。对于实施例公开的装置而言，由于其与实施例公开的方法相对应，所以描述的比较简单，相关之处参见方法部分说明即可。

对所公开的实施例的上述说明，使本领域专业技术人员能够实现或使用本发明。对这些实施例的多种修改对本领域的专业技术人员来说将是显而易见的，本文中所定义的一般原理可以在不脱离本发明的精神或范围的情况下，在其它实施例中实现。因此，本发明将不会被限制于本文所示的这些实施例，而是要符合与本文所公开的原理和新颖特点相一致的最宽的范围。