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

本发明专利技术属于GEMM运算加速领域，涉及一种GEMM运算加速器，包括主电路及与主电路相连接的从电路，其中：主电路针对输入的一批用于GEMM运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行GEMM运算，主电路合并从电路GEMM运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法API进行求解的传统方法得到运算结果后返回调用者。本发明专利技术的GEMM运算加速器利用动态分片，同时兼顾了线程级并行和指令级并行。本发明专利技术还提供一种基于GoogLeNet的图像处理加速方法。

【技术实现步骤摘要】
一种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的图像处理加速方法，包本文档来自技高网...

【技术保护点】
1.一种GEMM运算加速器，其特征在于，包括主电路及与主电路相连接的从电路，其中：/n主电路针对输入的一批用于GEMM运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行GEMM运算，主电路合并从电路GEMM运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法API进行求解的传统方法得到运算结果后返回调用者。/n

【技术特征摘要】
1.一种GEMM运算加速器，其特征在于，包括主电路及与主电路相连接的从电路，其中：
主电路针对输入的一批用于GEMM运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行GEMM运算，主电路合并从电路GEMM运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法API进行求解的传统方法得到运算结果后返回调用者。


2.根据权利要求1所述的GEMM运算加速器，其特征在于，动态分片过程包括：根据各个矩阵的规模以及使用的GPU架构和GPU相关参数从预先制定好的多个分片策略中选择在当前环境下的最优分片策略对矩阵进行分片。


3.根据权利要求2所述的GEMM运算加速器，其特征在于，预先制定好的多个分片策略应使分配给每一个矩阵片的workgroup大小一致。


4.根据权利要求3所述的GEMM运算加速器，其特征在于，使分配给每一个矩阵片的workgroup大小一致的方法包括：通过改变单个workgroup中每个workitem所负责运算的子片大小实现。


5.根据权利要求2所述的GEMM运算加速器，其特征在于，动态分片时采用一种平衡方法同时兼顾线程级并行和指令级并行。


6.根据权利要求5所述的GEMM运算加速器，其特征在于，平衡方法包括：
①、计算最优单个workgroup的workitem数量NWI：



其中：NMax_WG是单个workgroup最多所能包含的workitem的数量；NSIMD是单个CU所包含的SIMD的数量；
将NWI与预先制定好的多个分片策略中已有的单个workgroup的workitem数量参数进行比较，选择与NWI最接近的值：
min{abs(NWI-TWI_i)}
TWI_i为预先制定好的多个分片策略中单个wor...

【专利技术属性】
技术研发人员：羊志维，陆璐，
申请(专利权)人：华南理工大学，
类型：发明
国别省市：广东;44