一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法

本发明属于gemm运算加速领域，涉及一种gemm运算加速器及基于googlenet的图像处理加速方法。
背景技术：
：blas(basiclinearalgebrasubprogram，基础线性代数程序集)是一个api标准，用以规范发布基础线性代数操作的数值库(如向量或矩阵乘法)。最初发布于1979年，并用于建立更大的数值程序包(如lapack)，在高性能计算领域，blas被广泛使用。例如，linpack的运算成绩很大程度上取决于blas中子程序dgemm的表现。blas按照功能被分为三个级别：level1：向量-向量运算；level2：矩阵-向量运算；level3：矩阵-矩阵运算。而level3的blas包括gemm。gemm(generalmatrixmultiplication，通用矩阵乘法)是线性代数、机器学习、统计学和许多其他领域中的常见算法，形式为c＝α×a×b+β×c，a、b、c为矩阵，α、β为标量。由于矩阵乘法在各类科学应用中无处不在，所以gemm是blas优化的首要目标。gemm优化能在深度学习、天体物理学以及流体动力学等方面起到加速运算的作用。magma是新一代线性代数(la)gpu加速库的集合，由开发lapack和scalapack的团队设计并实现。magma适用于基于gpu的异构架构，它支持目前的la包和标准的接口，例如lapack和blas，以让相关开发研究人员能轻松地移植任何依赖la的软件组件。magma的主要优点在于，它可以使应用充分发挥当前多cpu(或多核cpu)和多gpu异构系统的威力，并在给定功耗限制下以最快的速度提供精确的解决方案。magma提供的加速库中包含名为vbatch的批量gemm运算的加速方案。rocm(radeonopencomputeplatform)是基于一系列开源项目的amdgpu计算生态，是首个面向hpc(highperformancecomputing)、超大规模gpu计算的开源软件开发平台。rocm为gpu计算带来了新的选择，即类unix、极简、模块化的软件开发。因为rocm生态系统由开源项目组成，所以它能一直保持活力，持续被优化以及扩展。开源项目包括机器学习框架(tensorflow、pytorch)、库(miopen、blas、rccl)、编程模型(hip)以及linuxkernel的支持等。rocm平台提供了hipblas_sgemm_batched一类api用于处理批量gemm运算，但是仅限于一批相同规模矩阵的gemm运算，而对于一批规模不定的矩阵，批量gemm运算的传统方法是循环执行hipblas_sgemm一类api，而magma作为目前与nvidia、amd合作的成熟的优化方案，相比于传统方法有所改进，它提供了magmablas_sgemm_vbatched一类api用于处理规模不定的矩阵的批量gemm运算。但在矩阵规模较小的情况下，gpu的利用率依旧很低，导致总的计算效率很低。例如googlenet，有57种卷积运算，而计算卷积的常用算法是将其转换为gemm(即c＝α×a×b+β×c的形式，a、b、c为矩阵，α、β为标量)再运算，对于转换后的矩阵，m(矩阵a和矩阵c的行数)、n(矩阵b和矩阵c的列数)、k(矩阵a的列数和矩阵b的行数)一般都小于1000，甚至有矩阵的m小于100，对于inception_3a/5x5_reduce中的卷积，转化为gemm后，其大小为m×n×k＝16×784×192，在mi50gpu上的性能不足峰值性能的1％，这是因为矩阵很小，分片后没有足够的workgroup来完全占据gpu。目前应用在cuda、rocm等平台下涉及规模不定的矩阵的批量gemm运算，只能是循环调用cublas_sgemm、hipblas_sgemm一类api去完成相关运算，由于在具体应用中涉及到的矩阵规模一般都较小(行列数小于等于1024)，导致gpu利用率很低，运算效率很差，而magma作为目前与nvidia、amd合作的成熟的优化方案，其中的vbatch方法相比于传统方法有所改进，它提供了magmablas_sgemm_vbatched一类api用于处理规模不定的矩阵的批量gemm运算。但在矩阵规模较小的情况下，gpu的利用率依旧很低，导致总的运算效率也很差。技术实现要素：针对矩阵规模较小的情况下，gemm运算效率低，gpu利用率低的问题，本发明提供一种gemm运算加速器。本发明还提供一种基于googlenet的图像处理加速方法。本发明的gemm运算加速器采用如下技术方案实现：一种gemm运算加速器，包括主电路及与主电路相连接的从电路，其中：主电路针对输入的一批用于gemm运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行gemm运算，主电路合并从电路gemm运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法api进行求解的传统方法得到运算结果后返回调用者。优选地，动态分片过程包括：根据各个矩阵的规模以及使用的gpu架构和gpu相关参数从预先制定好的多个分片策略中选择在当前环境下的最优分片策略对矩阵进行分片。优选地，预先制定好的多个分片策略应使分配给每一个矩阵片的workgroup大小一致。优选地，使分配给每一个矩阵片的workgroup大小一致的方法包括：通过改变单个workgroup中每个workitem所负责运算的子片大小实现。优选地，动态分片时采用一种平衡方法同时兼顾线程级并行和指令级并行。优选地，平衡方法包括：①、计算最优单个workgroup的workitem数量nwi：其中：nmax_wg是单个workgroup最多所能包含的workitem的数量；nsimd是单个cu所包含的simd的数量。将nwi与预先制定好的多个分片策略中已有的单个workgroup的workitem数量参数进行比较，选择与nwi最接近的值：min{abs(nwi-twi_i)}twi_i为预先制定好的多个分片策略中单个workgroup所包含的workitem数量。②、根据矩阵片大小小于输入矩阵大小的原则，筛选出可行分片策略；对可行分片策略分别进行计算得到对应的workgroup数量nwg_i：tm_i和tn_i是第i个分片策略的行数和列数，mj、nj为第j个gemm的矩阵c的行数和列数。③、选择与cu数的整数倍最接近的分片策略作为最优分片策略：min{nwg_imodncu}ncu为总的cu数量。优选地，平台包括cuda、rocm。本发明的基于googlenet的图像处理加速方法采用如下技术方案实现：一种基于googlenet的图像处理加速方法，包括：图像经过一系列的预处理后输入googlenet，经过若干层的处理后来到inception结构；将inception结构中的4个1×1卷积核所涉及的卷积运算转换成4个gemm运算输入gemm运算加速器中并行处理批量gemm运算；gemm运算加速器将运算结果返回googlenet；googlenet进行后续图像处理步骤。优选地，gemm运算加速器对于googlenet输入的gemm运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则进行动态分片，然后再对各个矩阵片并行进行gemm运算，合并gemm运算结果后将其返回给googlenet；若矩阵的行数或列数大于1024，则使用循环调用通用矩阵乘法api进行求解的传统方法得到运算结果后再将其返回给googlenet。本发明与现有技术相比，具有如下优点和有益效果：(1)相比于原生cuda、rocm平台下对于规模不等的矩阵的批量运算只能使用循环暴力求解的传统方法，以及新一代线性代数(la)gpu加速库magma的vbatch方法，本发明的gemm运算加速器利用动态分片，同时兼顾了线程级并行(tlp)和指令级并行(ilp)，最终使得批量gemm运算在矩阵行数和列数小于等于1024时耗时更短。(2)本发明gemm运算加速可基于rocm平台进行设计，封装成一个加速器供cuda、rocm、gpu等平台调用。同时，由于gemm矩阵乘法在各类科学应用中无处不在，本发明的gemm运算加速器可广泛应用于各类场景。例如在图像处理、深度学习、天体物理学以及流体动力学等场景。(3)本发明将googlenet的inception结构中涉及的4个1×1的卷积核转换成gemm后，不使用默认的运算方法，而是利用gemm运算加速器加速运算，达到减少googlenet在图像识别、图像分类等应用中的运算时间的目的。附图说明图1为一个实施例中gemm运算加速器工作流程图；图2为一个实施例中googlenetinception结构运算优化图。具体实施方式为使本发明的目的、技术方案和优点更加清楚明白，下面结合实例及附图对本发明作进一步详细的描述，但本发明的实施方式不限于此。实施例1一种gemm运算加速器，包括主电路及与主电路相连接的从电路，其中：主电路针对输入的一批用于gemm运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行gemm运算，主电路合并从电路gemm运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法api进行求解的传统方法得到运算结果后返回调用者。gemm运算加速器的工作流程如图1所示。下面以本发明提出的gemm运算加速器加速googlenetinception结构运算为例进行说明。googlenet是视觉领域竞赛ilsvrc2014年冠军模型(详见文献：szegedyc,liuw,jiay,etal.goingdeeperwithconvolutions[c]//proceedingsoftheieeeconferenceoncomputervisionandpatternrecognition.2015:1-9.)，其通过缩减参数的方法最大限度地节约计算资源，并首次提出inception结构，该结构利用多层感知器取代传统卷积神经网络中的广义线性结构，增加了网络宽度和深度，同时使用局部最优的稀疏结构取代原有卷积神经网络的全连接方式，最大限度地避免冗余。googlenet卷积神经网络由输入层、多层卷积层、多层子采样层和输出层构成，该结构共有22层，由于该神经网络层数很多，对于样本数据的抽象能力很强，同时参数个数很小，仅为5mb，有助于样本训练的同时能够快速收敛，并且该神经网络有3个loss值，可以进行不同层输出(详见文献：szegedyc,vanhouckev,ioffes,etal.rethinkingtheinceptionarchitectureforcomputervision[c]//proceedingsoftheieeeconferenceoncomputervisionandpatternrecognition.2016:2818-2826和文献leesg,sungy,kimyg,etal.variationsofalexnetandgooglenettoimprovekoreancharacterrecognitionperformance[j].journalofinformationprocessingsystems,2018,14(1))。对于googlenet的inception结构，如图2所示，inception结构中的卷积核包括4个1×1的卷积核、1个3×5的卷积核和1个5×5的卷积核。将inception结构中的4个1×1卷积核所涉及的卷积运算在转换成4个gemm运算后，默认是循环调用cuda、rocm等平台提供的通用矩阵乘法api进行求解，这种传统方法对gpu的利用率很低，导致运算效率不高。本发明将其中所涉及的矩阵输入到gemm运算加速器中，gemm运算加速器并行处理批量gemm运算，代替原本的求解方法，以加速googlenet在图像识别、图像分类等应用中的运算。针对googlenet输入的一批用于gemm运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024，若小于等于1024，则进行动态分片，即根据各个矩阵的规模以及使用的gpu的架构和相关参数从预先制定好的多个分片策略中选择当前环境下的最优策略对矩阵进行分片，然后再对各个矩阵片并行进行gemm运算，合并运算结果后将其返回给googlenet，若矩阵的行数或列数大于1024，则使用循环调用cuda、rocm等平台提供的通用矩阵乘法api进行求解的传统方法得到运算结果后再将其返回给googlenet。在一个优选的实施例中，制定的分片策略如表1所示：表1t_mt_nt_kworkitems/workgroup161681283232812864648128128648128641288128161682563232825664648256128648256641288256表1中：t_m、t_n、t_k分别为分片后的矩阵片的gemm(即c＝α×a×b+β×c的形式，a、b、c为分片后的矩阵片，α、β为标量)运算中矩阵片a和矩阵片c的行数、矩阵片b和矩阵片c的列数、矩阵片a的列数和矩阵片b的行数，workitems/workgroup代表单个workgroup中workitems的数量。相比于原生cuda、rocm平台下对于规模不等的矩阵的批量运算只能使用循环暴力求解的传统方法，以及新一代线性代数(la)gpu加速库magma的vbatch方法，本发明利用动态分片，同时兼顾了线程级并行(tlp)和指令级并行(ilp)，使得最终批量gemm运算性能在矩阵行数和列数小于等于1024时有较大地提升。对于预先设计好的一系列分片策略，为了避免不同规模的矩阵参与批量gemm运算所导致的线程空闲问题，预先制定好的一系列分片策略分片应该使分配给每一个矩阵片的workgroup大小一致，改变单个workgroup中的每个workitem所负责运算的子片大小即可。在一个实施例中，将分片策略根据矩阵片大小和单个workgroup包含workitem的数量，制定10种策略以供动态选择(如表1)，同时，还考虑不同大小的workgroup对子片大小的影响，例如对于矩阵片大小16×16、包含128个workitems的workgroup，子片大小为(16×16)/128＝2，所以令子片等于2×1。选择分片策略时需要兼顾线程级并行和指令级并行，当使用workitems较多的workgroup(例如包含256个workitems的workgroup)，这会提高线程级并行度，但是由于子片随之变小，导致指令级并行度降低，所以需要平衡两种并行度。在一个优选的实施例中，动态分片时采用一种平衡方法同时兼顾线程级并行和指令级并行。该平衡方法包括：①、针对每个gemm运算，先计算最优单个workgroup的workitem数量nwi：nmax_wg是单个workgroup最多所能包含的workitem的数量，nsimd是单个cu(即计算单元)所包含的simd(即矢量处理单元)的数量。通过与分片策略中已有的单个workgroup的workitem数量参数进行比较，选择最接近的值(已有的单个workgroup的workitem数量)。min{abs(nwi-twi_i)}twi_i为分片策略中单个workgroup所包含的workitem数量。②、然后根据矩阵片大小小于输入矩阵大小的原则，筛选出可行的分片策略。对筛选所得的分片策略分别进行计算得到对应的workgroup数量nwg_i：tm_i和tn_i是第i个分片策略的行数和列数，mj、nj为第j个gemm的矩阵c的行数和列数。③、求出对应分片策略的workgroup数量后，选择与cu数的整数倍最接近的分片策略作为最优分片策略：min{nwg_imodncu}ncu为总的cu数量。对于gemm运算，对矩阵c进行分片，分成若干大小为x×y的矩阵片，每个矩阵片是由矩阵a对应的行数据(x行)和矩阵b对应的列数据(y列)进行运算所得，而这些数据过多，gpu的单个vgpr(即向量通用寄存器)和lds(即本地数据缓存)无法一次性容纳，必须先将矩阵a的行数据按列分成若干矩阵片(x×k)，将矩阵b的列数据按行分成若干矩阵片(k×y)，然后将a矩阵片与b矩阵片的结果合并，才可得到最终结果。所有矩阵一开始存放在系统内存(systemmemory)中，先从系统内存中将对应的矩阵a的矩阵片和矩阵b的矩阵片放入gpu的lds中，然后从lds中将矩阵a和矩阵b的矩阵片根据分片策略中单个workgroup的workitem数量而分成若干子片，再将对应的子片放入vgpr中，计算这部分gemm运算，最后合并得出最终结果，以充分利用线程级并行性。具体如下：(一)选择分片策略如图1所示，googlenet输入一批规模大小不定的矩阵用以计算，由于本发明本质上是在gpu不能充分被利用的情况下尽可能增加gpu的利用率以提高运算效率，所以相比于传统方法和magmavbatch方法，本发明的gemm运算加速器更适用于小矩阵的矩阵运算，而googlenet在实际应用中经常涉及到行列数小于等于1024的矩阵运算，所以可以加速googlenet在图像识别、图像分类等场景的运算。gemm运算加速器对于googlenet输入的一批成对矩阵，先判断矩阵的行数和列数是否小于等于1024，对大于1024的矩阵，使用传统方法去计算，对小于等于1024的矩阵，使用优化后的方法计算。为了在预先制定好的分片策略中选出最优的分片策略，首先需要获取当前gpu架构以及相关参数，包括：单个workgroup最多所能包含的workitem的数量，单个cu所包含的simd的数量，总的cu数量。先计算最优单个workgroup的workitem数量nwi：nmax_wg是最大workgroup所能包含的workitem的数量，nsimd是单个cu所包含的simd的数量。通过与分片策略中已有的单个workgroup的workitem数量参数进行比较，选择最接近的值。min{abs(nwi-twi_i)}twi_i为分片策略中单个workgroup所包含的workitem数量。然后根据矩阵片大小小于输入矩阵大小的原则，筛选出可行的分片策略。对筛选所得的分片策略分别进行计算得到对应的workgroup数量nwg_i：tm_i和tn_i是第i个分片策略的行数和列数，mj、nj为第j个gemm的矩阵c的行数和列数。求出对应分片策略的workgroup数量后，选择与cu数的整数倍最接近的分片策略作为最优分片策略：min{nwg_imodncu}ncu为总的cu数量。(二)分片计算将存入系统内存中的成对矩阵按照选择的分片策略进行分片，为了提高传输效率，将对应的分好的矩阵片逐一存入各个cu(计算单元)中的lds(本地数据缓存)中，对lds中的矩阵a和矩阵b的矩阵片根据分片策略中单个workgroup的workitem数量分成若干子片，将对应的子片逐一传入各个simd(矢量处理单元)的vgpr(向量通用寄存器)中，再利用simd计算这部分gemm运算。(三)合并计算结果每个子片在运算结束后，将计算结果存回lds中原地址处，整个矩阵片在运算结束后，将计算结果存回系统内存中原地址处，在所有矩阵片都运算结束后，将整个gemm运算结果返回给googlenet。实施例2一种基于googlenet的图像处理加速方法，包括：图像在经过一系列的预处理后输入给googlenet，数据在经过若干层的处理后来到inception结构，inception结构中的卷积核包括4个1×1的卷积核、1个3×5的卷积核和1个5×5的卷积核，在cuda或rocm环境下，计算卷积的方法是将卷积转换为gemm(即c＝α×a×b+β×c的形式，a、b、c为矩阵，α、β为标量)再运算，对于转换后的矩阵，m(矩阵a和矩阵c的行数)、n(矩阵b和矩阵c的列数)、k(矩阵a的列数和矩阵b的行数)一般都小于1000，甚至有矩阵的m小于100，例如inception_3a/5x5_reduce中的卷积，转化为gemm后，其大小为m×n×k＝16×784×192，在mi50gpu上的性能不足峰值性能的1％，这是因为矩阵很小，分片后没有足够的workgroup来完全占据gpu。在cuda或rocm环境下，将会串行调用gemm子例程对inception结构所涉及到的矩阵运算进行求解，而本发明所涉及到的gemm运算加速器将会对4个1×1的卷积核所涉及到的矩阵运算放在一起进行统一处理，以提高gpu的利用率，最终达到图像处理加速的效果。上述实例为本发明较佳的实施方式，对本发明的目的、技术方案和优点进行了进一步详细说明，但本发明的实施方式并不受上述实例的限制，其他的任何未背离本发明的精神实质与原理下所作的任何修改、等同替换、改进等，均应为等效的置换方式，都包含在本发明的保护范围之内。当前第1页12