CN110246078A - 一种基于嵌入式gpu和卷积计算的图像处理方法和装置 - Google Patents

一种基于嵌入式gpu和卷积计算的图像处理方法和装置 Download PDF

Info

Publication number
CN110246078A
CN110246078A CN201910472732.3A CN201910472732A CN110246078A CN 110246078 A CN110246078 A CN 110246078A CN 201910472732 A CN201910472732 A CN 201910472732A CN 110246078 A CN110246078 A CN 110246078A
Authority
CN
China
Prior art keywords
matrix
convolution kernel
submatrix
convolution
convolutional calculation
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Granted
Application number
CN201910472732.3A
Other languages
English (en)
Other versions
CN110246078B (zh
Inventor
姜宏旭
王玺坤
李波
张永华
林珂玉
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Hangzhou Innovation Research Institute of Beihang University
Original Assignee
Beihang University
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Beihang University filed Critical Beihang University
Priority to CN201910472732.3A priority Critical patent/CN110246078B/zh
Publication of CN110246078A publication Critical patent/CN110246078A/zh
Application granted granted Critical
Publication of CN110246078B publication Critical patent/CN110246078B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
    • G06F17/10—Complex mathematical operations
    • G06F17/15—Correlation function computation including computation of convolution operations
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F17/00—Digital computing or data processing equipment or methods, specially adapted for specific functions
    • G06F17/10—Complex mathematical operations
    • G06F17/16—Matrix or vector computation, e.g. matrix-matrix or matrix-vector multiplication, matrix factorization
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/20—Processor architectures; Processor configuration, e.g. pipelining

Landscapes

  • Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Mathematical Physics (AREA)
  • Theoretical Computer Science (AREA)
  • Pure & Applied Mathematics (AREA)
  • Mathematical Optimization (AREA)
  • Mathematical Analysis (AREA)
  • Data Mining & Analysis (AREA)
  • Computational Mathematics (AREA)
  • Computing Systems (AREA)
  • Algebra (AREA)
  • Databases & Information Systems (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Complex Calculations (AREA)

Abstract

本发明公开了一种基于嵌入式GPU和卷积计算的图像处理方法和装置，针对SSD算法中的卷积计算进行优化，采用内存优化的卷积扩展对输入图像进行矩阵变换，利用CUDA并行处理形成中间矩阵，同时采用卷积核矩阵行列扩充对齐，卷积核矩阵扩充后做分块处理以减少运算时内存开销，最后采用CUDA库中高度优化的cuBLAS矩阵乘法函数进行卷积计算并行加速，最后合并输出矩阵。本发明提供的方法能够降低内存开销，提高算法的性能，同时发挥GPU并行控制的优势，减少矩阵乘法时间，提高计算效率。

Description

一种基于嵌入式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平台Jetson TX2上，经大量实验证明，本发明提供的方法平均内存使用效率提升了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平台NVIDIA Jetson TX2上进行，以GPU的时钟频率为基准，对程序的运行时间和内存进行了测量，每个算法运行10次，并取平均值，实验结果如7所示。
通过实验结果可知，优化后的卷积计算在运行时间和运行内存上，与优化前相比性能有了大幅度提升，内存开销大大减少，运行时间得到提高。特别对于SSD算法的卷积层，3×3卷积，图像输入分辨率在300×300的情况下，运算时间提高了90％以上，运行内存减小了45％以上。
本说明书中各个实施例采用递进的方式描述，每个实施例重点说明的都是与其他实施例的不同之处，各个实施例之间相同相似部分互相参见即可。对于实施例公开的装置而言，由于其与实施例公开的方法相对应，所以描述的比较简单，相关之处参见方法部分说明即可。
对所公开的实施例的上述说明，使本领域专业技术人员能够实现或使用本发明。对这些实施例的多种修改对本领域的专业技术人员来说将是显而易见的，本文中所定义的一般原理可以在不脱离本发明的精神或范围的情况下，在其它实施例中实现。因此，本发明将不会被限制于本文所示的这些实施例，而是要符合与本文所公开的原理和新颖特点相一致的最宽的范围。

Claims (8)

1.一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，包括：
S1：采用内存优化的卷积扩展方法对输入图像进行矩阵变换和CUDA并行处理，得到中间矩阵；
S2：对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；
S3：对中间矩阵和卷积核子矩阵进行卷积计算；
S4：对卷积计算后输出的子矩阵进行合并，得到输出矩阵。
2.根据权利要求1所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，在步骤S1中，采用内存优化的卷积扩展方法对输入图像进行矩阵变换的具体步骤包括：
S11：在输入图像对应的输入矩阵I[0:ih,0:iw]上，选取[0:ih,0:sw]为滑动窗口，滑动窗口大小为ih*sw；ih为输入矩阵的行，iw为输入矩阵的列；
S12：滑动窗口[0:ih,0:sw]依次在输入矩阵I[0:in,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为中间矩阵的列。
3.根据权利要求1所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，步骤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为分块后卷积核第二子矩阵的列。
4.根据权利要求1所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，在步骤S3中，利用CUDA库中高度优化的cuBLAS函数对中间矩阵和卷积核子矩阵进行卷积计算。
5.根据权利要求4所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，步骤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库对象。
6.根据权利要求5所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，步骤S4具体包括：
将第一子矩阵O1[0:Oh,0:Ow1]和第二子矩阵O2[0:Oh,0:Ow2]合并为输出矩阵O[0:Oh,0:Ow]，其中，Oh为输出矩阵的行，Ow为输出矩阵的列。
7.根据权利要求1～6任意一项所述的一种基于嵌入式GPU和卷积计算的图像处理方法，其特征在于，该方法由CPU和GPU协同处理实现，其中，GPU负责执行步骤S1和步骤S3，CPU负责执行步骤S2和步骤S4。
8.一种基于嵌入式GPU和卷积计算的图像处理装置，其特征在于，包括：
卷积扩展模块，用于采用内存优化的卷积扩展方法对输入图像进行矩阵变换和CUDA并行处理，得到中间矩阵；
扩充分块处理模块，用于对输入图像进行卷积核矩阵行列扩充，得到卷积核临时矩阵，对所述卷积核临时矩阵进行分块处理，生成卷积核子矩阵；
卷积计算模块，用于对中间矩阵和卷积核子矩阵进行卷积计算；
合并模块，用于对卷积计算后输出子矩阵进行合并，得到输出矩阵。
CN201910472732.3A 2019-05-31 2019-05-31 一种基于嵌入式gpu和卷积计算的图像处理方法和装置 Active CN110246078B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN201910472732.3A CN110246078B (zh) 2019-05-31 2019-05-31 一种基于嵌入式gpu和卷积计算的图像处理方法和装置

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN201910472732.3A CN110246078B (zh) 2019-05-31 2019-05-31 一种基于嵌入式gpu和卷积计算的图像处理方法和装置

Publications (2)

Publication Number Publication Date
CN110246078A true CN110246078A (zh) 2019-09-17
CN110246078B CN110246078B (zh) 2020-11-03

Family

ID=67885759

Family Applications (1)

Application Number Title Priority Date Filing Date
CN201910472732.3A Active CN110246078B (zh) 2019-05-31 2019-05-31 一种基于嵌入式gpu和卷积计算的图像处理方法和装置

Country Status (1)

Country Link
CN (1) CN110246078B (zh)

Cited By (6)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111381968A (zh) * 2020-03-11 2020-07-07 中山大学 一种高效运行深度学习任务的卷积运算优化方法及系统
CN112561943A (zh) * 2020-12-23 2021-03-26 清华大学 一种基于脉动阵列卷积运算数据复用的图像处理方法
CN112991142A (zh) * 2021-03-31 2021-06-18 腾讯科技（深圳）有限公司 图像数据的矩阵运算方法、装置、设备及存储介质
CN113240570A (zh) * 2021-04-13 2021-08-10 华南理工大学 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法
CN113536220A (zh) * 2020-04-21 2021-10-22 中科寒武纪科技股份有限公司 运算方法、处理器及相关产品
CN115985465A (zh) * 2023-03-21 2023-04-18 天津医科大学总医院 基于时序的肌电信号特征提取方法、装置、设备及存储介质

Citations (7)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20130080073A1 (en) * 2010-06-11 2013-03-28 Waters Technologies Corporation Techniques for mass spectrometry peak list computation using parallel processing
CN104299216A (zh) * 2014-10-22 2015-01-21 北京航空航天大学 基于多尺度各向异性分解和低秩分析的多模态医学图像融合方法
CN108631925A (zh) * 2017-03-24 2018-10-09 中兴通讯股份有限公司 一种准循环低密度奇偶校验编码处理方法及装置
CN108959794A (zh) * 2018-07-13 2018-12-07 北京航空航天大学 一种基于深度学习的结构频响动力学模型修正方法
CN109325589A (zh) * 2017-07-31 2019-02-12 华为技术有限公司 卷积计算方法及装置
CN109389059A (zh) * 2018-09-26 2019-02-26 华南理工大学 一种基于cnn-lstm网络的p300检测方法
CN109814986A (zh) * 2017-11-20 2019-05-28 上海寒武纪信息科技有限公司 任务并行处理方法、存储介质、计算机设备、装置和系统

Patent Citations (7)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20130080073A1 (en) * 2010-06-11 2013-03-28 Waters Technologies Corporation Techniques for mass spectrometry peak list computation using parallel processing
CN104299216A (zh) * 2014-10-22 2015-01-21 北京航空航天大学 基于多尺度各向异性分解和低秩分析的多模态医学图像融合方法
CN108631925A (zh) * 2017-03-24 2018-10-09 中兴通讯股份有限公司 一种准循环低密度奇偶校验编码处理方法及装置
CN109325589A (zh) * 2017-07-31 2019-02-12 华为技术有限公司 卷积计算方法及装置
CN109814986A (zh) * 2017-11-20 2019-05-28 上海寒武纪信息科技有限公司 任务并行处理方法、存储介质、计算机设备、装置和系统
CN108959794A (zh) * 2018-07-13 2018-12-07 北京航空航天大学 一种基于深度学习的结构频响动力学模型修正方法
CN109389059A (zh) * 2018-09-26 2019-02-26 华南理工大学 一种基于cnn-lstm网络的p300检测方法

Non-Patent Citations (2)

* Cited by examiner, † Cited by third party
Title
VICTOR PODLOZHNYUK: "NVIDIA: Image Convolution with CUDA", 《HTTP://WWW.DOC88.COM/P-1754504897085.HTML》 *
姜宏旭，等: "基于FPGA的高效机载视频采集及预处理方法", 《北京航空航天大学学报》 *

Cited By (11)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN111381968A (zh) * 2020-03-11 2020-07-07 中山大学 一种高效运行深度学习任务的卷积运算优化方法及系统
CN111381968B (zh) * 2020-03-11 2023-04-25 中山大学 一种高效运行深度学习任务的卷积运算优化方法及系统
CN113536220A (zh) * 2020-04-21 2021-10-22 中科寒武纪科技股份有限公司 运算方法、处理器及相关产品
CN112561943A (zh) * 2020-12-23 2021-03-26 清华大学 一种基于脉动阵列卷积运算数据复用的图像处理方法
CN112561943B (zh) * 2020-12-23 2022-11-22 清华大学 一种基于脉动阵列卷积运算数据复用的图像处理方法
CN112991142A (zh) * 2021-03-31 2021-06-18 腾讯科技（深圳）有限公司 图像数据的矩阵运算方法、装置、设备及存储介质
WO2022206556A1 (zh) * 2021-03-31 2022-10-06 腾讯科技（深圳）有限公司 图像数据的矩阵运算方法、装置、设备及存储介质
CN112991142B (zh) * 2021-03-31 2023-06-16 腾讯科技（深圳）有限公司 图像数据的矩阵运算方法、装置、设备及存储介质
CN113240570A (zh) * 2021-04-13 2021-08-10 华南理工大学 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法
CN115985465A (zh) * 2023-03-21 2023-04-18 天津医科大学总医院 基于时序的肌电信号特征提取方法、装置、设备及存储介质
CN115985465B (zh) * 2023-03-21 2023-07-07 天津医科大学总医院 基于时序的肌电信号特征提取方法、装置、设备及存储介质

Also Published As

Publication number Publication date
CN110246078B (zh) 2020-11-03

Similar Documents

Publication Publication Date Title
CN110246078A (zh) 一种基于嵌入式gpu和卷积计算的图像处理方法和装置
CN104915322B (zh) 一种卷积神经网络硬件加速方法
JP6977239B2 (ja) 行列乗算器
Wang et al. Sparse-YOLO: Hardware/software co-design of an FPGA accelerator for YOLOv2
JP6771018B2 (ja) 二次元配列プロセッサの性能向上
CN107341544A (zh) 一种基于可分割阵列的可重构加速器及其实现方法
CN109086244A (zh) 一种基于向量处理器的矩阵卷积向量化实现方法
CN101398753A (zh) 用于执行扫描运算的系统、方法及计算机程序产品
TWI827792B (zh) 多路徑神經網路、資源配置的方法及多路徑神經網路分析器
CN103049241A (zh) 一种提高cpu+gpu异构装置计算性能的方法
CN105739951A (zh) 一种基于gpu的l1最小化问题快速求解方法
CN104699464A (zh) 一种基于依赖网格的指令级并行调度方法
Motamedi et al. Fast and energy-efficient CNN inference on IoT devices
Chen et al. Phonebit: Efficient gpu-accelerated binary neural network inference engine for mobile phones
CN110377874B (zh) 卷积运算方法及系统
Clarke et al. Fupermod: A framework for optimal data partitioning for parallel scientific applications on dedicated heterogeneous hpc platforms
CN110716751B (zh) 高并行度计算平台、系统及计算实现方法
CN113313252A (zh) 一种基于脉动阵列的深度可分离卷积实现方法
Mahmoudi et al. Efficient exploitation of heterogeneous platforms for images features extraction
CN115293978A (zh) 卷积运算电路和方法、图像处理设备
US20090326888A1 (en) Vectorized parallel collision detection pipeline
Zhu et al. Parallelized force-directed edge bundling on the GPU
CN101930079A (zh) 一种地震勘探相关/叠加数据处理方法
JP4950325B2 (ja) モンテカルロ法の効率的な並列処理手法
Sredojevic et al. Structured deep neural network pruning via matrix pivoting

Legal Events

Date Code Title Description
PB01 Publication
PB01 Publication
SE01 Entry into force of request for substantive examination
SE01 Entry into force of request for substantive examination
GR01 Patent grant
GR01 Patent grant
TR01 Transfer of patent right
TR01 Transfer of patent right

Effective date of registration: 20210428

Address after: 310000 No. 18 Chuanghui Street, Changhe Street, Binjiang District, Hangzhou City, Zhejiang Province

Patentee after: BUAA HANGZHOU INNOVATION INSTITUTE

Address before: 100191 Haidian District, Xueyuan Road, No. 37,

Patentee before: BEIHANG University