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

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

Info

Publication number
CN113240570B
CN113240570B CN202110392571.4A CN202110392571A CN113240570B CN 113240570 B CN113240570 B CN 113240570B CN 202110392571 A CN202110392571 A CN 202110392571A CN 113240570 B CN113240570 B CN 113240570B
Authority
CN
China
Prior art keywords
matrix
gemm
accelerator
fragmentation
gemm operation
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.)
Active
Application number
CN202110392571.4A
Other languages
English (en)
Other versions
CN113240570A (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.)
Shenzhen Aitesi Information Technology Co ltd
Original Assignee
South China University of Technology SCUT
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 South China University of Technology SCUT filed Critical South China University of Technology SCUT
Priority to CN202110392571.4A priority Critical patent/CN113240570B/zh
Publication of CN113240570A publication Critical patent/CN113240570A/zh
Application granted granted Critical
Publication of CN113240570B publication Critical patent/CN113240570B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • 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
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5027—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06N—COMPUTING ARRANGEMENTS BASED ON SPECIFIC COMPUTATIONAL MODELS
    • G06N3/00—Computing arrangements based on biological models
    • G06N3/02—Neural networks
    • G06N3/04—Architecture, e.g. interconnection topology
    • G06N3/045—Combinations of networks

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • Software Systems (AREA)
  • General Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • Data Mining & Analysis (AREA)
  • Computational Linguistics (AREA)
  • Biophysics (AREA)
  • Evolutionary Computation (AREA)
  • General Health & Medical Sciences (AREA)
  • Molecular Biology (AREA)
  • Computing Systems (AREA)
  • Biomedical Technology (AREA)
  • Artificial Intelligence (AREA)
  • Mathematical Physics (AREA)
  • Life Sciences & Earth Sciences (AREA)
  • Health & Medical Sciences (AREA)
  • Management, Administration, Business Operations System, And Electronic Commerce (AREA)

Abstract

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

Description

一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法
技术领域
本发明属于GEMM运算加速领域，涉及一种GEMM运算加速器及基于 GoogLeNet的图像处理加速方法。
背景技术
BLAS(Basic Linear Algebra Subprogram，基础线性代数程序集)是一个API 标准，用以规范发布基础线性代数操作的数值库(如向量或矩阵乘法)。最初发 布于1979年，并用于建立更大的数值程序包(如LAPACK)，在高性能计算领 域，BLAS被广泛使用。例如，LINPACK的运算成绩很大程度上取决于BLAS 中子程序DGEMM的表现。BLAS按照功能被分为三个级别：Level 1：向量-向 量运算；Level 2：矩阵-向量运算；Level 3：矩阵-矩阵运算。而Level 3的BLAS 包括GEMM。
GEMM(General Matrix Multiplication，通用矩阵乘法)是线性代数、机器 学习、统计学和许多其他领域中的常见算法，形式为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(Radeon Open Compute platform)是基于一系列开源项目的AMD GPU计算生态，是首个面向HPC(High Performance Computing)、超大规模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，在MI50 GPU上的性能不足峰值性 能的1％，这是因为矩阵很小，分片后没有足够的work group来完全占据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相关参数从预先制定好的多个分片策略中选择在当前环境下的最优分片策 略对矩阵进行分片。
优选地，预先制定好的多个分片策略应使分配给每一个矩阵片的work group 大小一致。
优选地，使分配给每一个矩阵片的work group大小一致的方法包括：通过改 变单个work group中每个work item所负责运算的子片大小实现。
优选地，动态分片时采用一种平衡方法同时兼顾线程级并行和指令级并行。
优选地，平衡方法包括：
①、计算最优单个work group的work item数量NWI：
Figure BDA0003017311850000041
其中：NMax_WG是单个work group最多所能包含的work item的数量；NSIMD是 单个CU所包含的SIMD的数量。
将NWI与预先制定好的多个分片策略中已有的单个work group的work item数 量参数进行比较，选择与NWI最接近的值：
min{abs(NWI-TWI_i)}
TWI_i为预先制定好的多个分片策略中单个work group所包含的work item数 量。
②、根据矩阵片大小小于输入矩阵大小的原则，筛选出可行分片策略；对 可行分片策略分别进行计算得到对应的work group数量NWG_i：
Figure BDA0003017311850000043
TM_i和TN_i是第i个分片策略的行数和列数，Mj、Nj为第j个GEMM的矩阵C的 行数和列数。
③、选择与CU数的整数倍最接近的分片策略作为最优分片策略：
min{NWG_imod NCU}
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为一个实施例中GoogLeNet inception结构运算优化图。
具体实施方式
为使本发明的目的、技术方案和优点更加清楚明白，下面结合实例及附图 对本发明作进一步详细的描述，但本发明的实施方式不限于此。
实施例1
一种GEMM运算加速器，包括主电路及与主电路相连接的从电路，其中： 主电路针对输入的一批用于GEMM运算的规模不等的矩阵，先判断矩阵的行数 和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从 电路对各个矩阵片进行GEMM运算，主电路合并从电路GEMM运算结果后返回 调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵 乘法API进行求解的传统方法得到运算结果后返回调用者。GEMM运算加速器的 工作流程如图1所示。
下面以本发明提出的GEMM运算加速器加速GoogLeNet inception结构运 算为例进行说明。
GoogLeNet是视觉领域竞赛ILSVRC 2014年冠军模型(详见文献：Szegedy C,LiuW,Jia Y,et al.Going deeper with convolutions[C]//Proceedings of the IEEEconference on computer vision and pattern recognition.2015:1-9.)，其通过缩减参数的方法最大限度地节约计算资源，并首次提出inception结构，该结构利用多 层感知器取代传统卷积神经网络中的广义线性结构，增加了网络宽度和深度， 同时使用局部最优的稀疏结构取代原有卷积神经网络的全连接方式，最大限度 地避免冗余。GoogLeNet卷积神经网络由输入层、多层卷积层、多层子采样层 和输出层构成，该结构共有22层，由于该神经网络层数很多，对于样本数据的 抽象能力很强，同时参数个数很小，仅为5MB，有助于样本训练的同时能够快 速收敛，并且该神经网络有3个loss值，可以进行不同层输出(详见文献：Szegedy C,Vanhoucke V,Ioffe S,et al.Rethinking the inception architecturefor computer vision[C]//Proceedings of the IEEE conference on computer visionand pattern recognition.2016:2818-2826和文献Lee S G,Sung Y,Kim Y G,etal.Variations of AlexNet and GoogLeNet to Improve Korean CharacterRecognition Performance[J]. Journal of Information Processing Systems,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所示：
表1
T_M T_N T_K Work Items/Work Group
16 16 8 128
32 32 8 128
64 64 8 128
128 64 8 128
64 128 8 128
16 16 8 256
32 32 8 256
64 64 8 256
128 64 8 256
64 128 8 256
表1中：T_M、T_N、T_K分别为分片后的矩阵片的GEMM(即C＝α×A×B+ β×C的形式，A、B、C为分片后的矩阵片，α、β为标量)运算中矩阵片A和矩 阵片C的行数、矩阵片B和矩阵片C的列数、矩阵片A的列数和矩阵片B的行数， Work Items/Work Group代表单个Work Group中Work Items的数量。
相比于原生CUDA、ROCm平台下对于规模不等的矩阵的批量运算只能使用 循环暴力求解的传统方法，以及新一代线性代数(LA)GPU加速库MAGMA的 vbatch方法，本发明利用动态分片，同时兼顾了线程级并行(TLP)和指令级并 行(ILP)，使得最终批量GEMM运算性能在矩阵行数和列数小于等于1024时有 较大地提升。
对于预先设计好的一系列分片策略，为了避免不同规模的矩阵参与批量 GEMM运算所导致的线程空闲问题，预先制定好的一系列分片策略分片应该使 分配给每一个矩阵片的work group大小一致，改变单个work group中的每个work item所负责运算的子片大小即可。
在一个实施例中，将分片策略根据矩阵片大小和单个work group包含work item的数量，制定10种策略以供动态选择(如表1)，同时，还考虑不同大小的 work group对子片大小的影响，例如对于矩阵片大小16×16、包含128个work items的work group，子片大小为(16×16)/128＝2，所以令子片等于2×1。
选择分片策略时需要兼顾线程级并行和指令级并行，当使用work items较多 的work group(例如包含256个work items的work group)，这会提高线程级并行度， 但是由于子片随之变小，导致指令级并行度降低，所以需要平衡两种并行度。
在一个优选的实施例中，动态分片时采用一种平衡方法同时兼顾线程级并 行和指令级并行。该平衡方法包括：
①、针对每个GEMM运算，先计算最优单个work group的work item数量NWI：
Figure BDA0003017311850000091
NMax_WG是单个work group最多所能包含的work item的数量，NSIMD是单个 CU(即计算单元)所包含的SIMD(即矢量处理单元)的数量。
通过与分片策略中已有的单个work group的work item数量参数进行比较，选 择最接近的值(已有的单个work group的work item数量)。
min{abs(NWI-TWI_i)}
TWI_i为分片策略中单个work group所包含的work item数量。
②、然后根据矩阵片大小小于输入矩阵大小的原则，筛选出可行的分片策 略。对筛选所得的分片策略分别进行计算得到对应的work group数量NWG_i：
Figure BDA0003017311850000093
TM_i和TN_i是第i个分片策略的行数和列数，Mj、Nj为第j个GEMM的矩阵C的 行数和列数。
③、求出对应分片策略的work group数量后，选择与CU数的整数倍最接近 的分片策略作为最优分片策略：
min{NWG_i mod NCU}
NCU为总的CU数量。
对于GEMM运算，对矩阵C进行分片，分成若干大小为X×Y的矩阵片，每 个矩阵片是由矩阵A对应的行数据(X行)和矩阵B对应的列数据(Y列)进行运 算所得，而这些数据过多，GPU的单个VGPR(即向量通用寄存器)和LDS(即 本地数据缓存)无法一次性容纳，必须先将矩阵A的行数据按列分成若干矩阵片 (X×K)，将矩阵B的列数据按行分成若干矩阵片(K×Y)，然后将A矩阵片与B 矩阵片的结果合并，才可得到最终结果。
所有矩阵一开始存放在***内存(System Memory)中，先从***内存中将 对应的矩阵A的矩阵片和矩阵B的矩阵片放入GPU的LDS中，然后从LDS中 将矩阵A和矩阵B的矩阵片根据分片策略中单个work group的work item数量 而分成若干子片，再将对应的子片放入VGPR中，计算这部分GEMM运算，最 后合并得出最终结果，以充分利用线程级并行性。
具体如下：
(一)选择分片策略
如图1所示，GoogLeNet输入一批规模大小不定的矩阵用以计算，由于本发 明本质上是在GPU不能充分被利用的情况下尽可能增加GPU的利用率以提高运 算效率，所以相比于传统方法和MAGMA vbatch方法，本发明的GEMM运算加速 器更适用于小矩阵的矩阵运算，而GoogLeNet在实际应用中经常涉及到行列数小 于等于1024的矩阵运算，所以可以加速GoogLeNet在图像识别、图像分类等场景 的运算。GEMM运算加速器对于GoogLeNet输入的一批成对矩阵，先判断矩阵的 行数和列数是否小于等于1024，对大于1024的矩阵，使用传统方法去计算，对 小于等于1024的矩阵，使用优化后的方法计算。
为了在预先制定好的分片策略中选出最优的分片策略，首先需要获取当前 GPU架构以及相关参数，包括：单个work group最多所能包含的work item的数量， 单个CU所包含的SIMD的数量，总的CU数量。
先计算最优单个work group的work item数量NWI：
Figure BDA0003017311850000111
NMax_WG是最大work group所能包含的work item的数量，NSIMD是单个CU所 包含的SIMD的数量。通过与分片策略中已有的单个work group的work item数量 参数进行比较，选择最接近的值。
min{abs(NWI-TWI_i)}
TWI_i为分片策略中单个work group所包含的work item数量。
然后根据矩阵片大小小于输入矩阵大小的原则，筛选出可行的分片策略。 对筛选所得的分片策略分别进行计算得到对应的work group数量NWG_i：
Figure BDA0003017311850000113
TM_i和TN_i是第i个分片策略的行数和列数，Mj、Nj为第j个GEMM的矩阵C的 行数和列数。
求出对应分片策略的work group数量后，选择与CU数的整数倍最接近的分 片策略作为最优分片策略：
min{NWG_i mod NCU}
NCU为总的CU数量。
(二)分片计算
将存入***内存中的成对矩阵按照选择的分片策略进行分片，为了提高传 输效率，将对应的分好的矩阵片逐一存入各个CU(计算单元)中的LDS(本地 数据缓存)中，对LDS中的矩阵A和矩阵B的矩阵片根据分片策略中单个work group的work item数量分成若干子片，将对应的子片逐一传入各个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，在MI50 GPU上的性能 不足峰值性能的1％，这是因为矩阵很小，分片后没有足够的work group来完全 占据GPU。在CUDA或ROCm环境下，将会串行调用GEMM子例程对inception 结构所涉及到的矩阵运算进行求解，而本发明所涉及到的GEMM运算加速器将 会对4个1×1的卷积核所涉及到的矩阵运算放在一起进行统一处理，以提高 GPU的利用率，最终达到图像处理加速的效果。
上述实例为本发明较佳的实施方式，对本发明的目的、技术方案和优点进 行了进一步详细说明，但本发明的实施方式并不受上述实例的限制，其他的任 何未背离本发明的精神实质与原理下所作的任何修改、等同替换、改进等，均 应为等效的置换方式，都包含在本发明的保护范围之内。

Claims (8)

1.一种GEMM运算加速器，其特征在于，包括主电路及与主电路相连接的从电路，其中：
主电路针对输入的一批用于GEMM运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则对矩阵进行动态分片，然后从电路对各个矩阵片进行GEMM运算，主电路合并从电路GEMM运算结果后返回调用者；若矩阵的行数或列数大于1024，则使用循环调用平台提供的通用矩阵乘法API进行求解的传统方法得到运算结果后返回调用者；
动态分片时采用一种平衡方法同时兼顾线程级并行和指令级并行，平衡方法包括：
①、计算最优单个work group的work item数量NWI：
Figure FDA0003766900560000011
其中：NMax_WG是单个work group最多所能包含的work item的数量；NSIMD是单个CU所包含的SIMD的数量；
将NWI与预先制定好的多个分片策略中已有的单个work group的work item数量参数进行比较，选择与NWI最接近的值：
min{abs(NWI-TWI_i)}
TWI_i为预先制定好的多个分片策略中单个work group所包含的work item数量；
②、根据矩阵片大小小于输入矩阵大小的原则，筛选出可行分片策略；对可行分片策略分别进行计算得到对应的work group数量NWG_i：
Figure FDA0003766900560000012
TM_i和TN_i是第i个分片策略的行数和列数，Mj、Nj为第j个GEMM的矩阵C的行数和列数；
③、选择与CU数的整数倍最接近的分片策略作为最优分片策略：
min{NWG_imod NCU}
NCU为总的CU数量。
2.根据权利要求1所述的GEMM运算加速器，其特征在于，动态分片过程包括：根据各个矩阵的规模以及使用的GPU架构和GPU相关参数从预先制定好的多个分片策略中选择在当前环境下的最优分片策略对矩阵进行分片。
3.根据权利要求2所述的GEMM运算加速器，其特征在于，预先制定好的多个分片策略应使分配给每一个矩阵片的work group大小一致。
4.根据权利要求3所述的GEMM运算加速器，其特征在于，使分配给每一个矩阵片的workgroup大小一致的方法包括：通过改变单个work group中每个work item所负责运算的子片大小实现。
5.根据权利要求1所述的GEMM运算加速器，其特征在于，平台包括CUDA、ROCm。
6.一种基于GoogLeNet的图像处理加速方法，其特征在于，基于权利要求1-4中任一所述的GEMM运算加速器实现，包括：
图像经过一系列的预处理后输入GoogLeNet，经过若干层的处理后来到inception结构；
将inception结构中的4个1×1卷积核所涉及的卷积运算转换成4个GEMM运算输入GEMM运算加速器中并行处理批量GEMM运算；
GEMM运算加速器将运算结果返回GoogLeNet；
GoogLeNet进行后续图像处理步骤。
7.根据权利要求6所述的图像处理加速方法，其特征在于，GEMM运算加速器对于GoogLeNet输入的GEMM运算的规模不等的矩阵，先判断矩阵的行数和列数是否小于等于1024：若小于等于1024，则进行动态分片，然后再对各个矩阵片并行进行GEMM运算，合并GEMM运算结果后将其返回给GoogLeNet；若矩阵的行数或列数大于1024，则使用循环调用通用矩阵乘法API进行求解的传统方法得到运算结果后再将其返回给GoogLeNet。
8.根据权利要求7所述的图像处理加速方法，其特征在于，动态分片过程包括：根据各个矩阵的规模以及使用的GPU架构和GPU相关参数从预先制定好的多个分片策略中选择在当前环境下的最优分片策略对矩阵进行分片。
CN202110392571.4A 2021-04-13 2021-04-13 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法 Active CN113240570B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202110392571.4A CN113240570B (zh) 2021-04-13 2021-04-13 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202110392571.4A CN113240570B (zh) 2021-04-13 2021-04-13 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法

Publications (2)

Publication Number Publication Date
CN113240570A CN113240570A (zh) 2021-08-10
CN113240570B true CN113240570B (zh) 2023-01-06

Family

ID=77128680

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202110392571.4A Active CN113240570B (zh) 2021-04-13 2021-04-13 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法

Country Status (1)

Country Link
CN (1) CN113240570B (zh)

Citations (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN110245751A (zh) * 2017-08-31 2019-09-17 北京中科寒武纪科技有限公司 一种gemm运算运算方法及装置

Family Cites Families (5)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN105808309B (zh) * 2016-03-08 2019-04-05 中国科学院软件研究所 一种基于申威平台的基础线性代数库blas三级函数gemm的高性能实现方法
US10073815B2 (en) * 2016-05-31 2018-09-11 Palo Alto Research Cener Incorporated System and method for speeding up general matrix-matrix multiplication on the GPU
US10657442B2 (en) * 2018-04-19 2020-05-19 International Business Machines Corporation Deep learning accelerator architecture with chunking GEMM
US11580386B2 (en) * 2019-03-18 2023-02-14 Electronics And Telecommunications Research Institute Convolutional layer acceleration unit, embedded system having the same, and method for operating the embedded system
CN110246078B (zh) * 2019-05-31 2020-11-03 北京航空航天大学 一种基于嵌入式gpu和卷积计算的图像处理方法和装置

Patent Citations (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN110245751A (zh) * 2017-08-31 2019-09-17 北京中科寒武纪科技有限公司 一种gemm运算运算方法及装置

Also Published As

Publication number Publication date
CN113240570A (zh) 2021-08-10

Similar Documents

Publication Publication Date Title
Jouppi et al. A domain-specific supercomputer for training deep neural networks
Guo et al. Software-hardware codesign for efficient neural network acceleration
Norrie et al. The design process for Google's training chips: TPUv2 and TPUv3
CN109543830B (zh) 一种用于卷积神经网络加速器的拆分累加器
Lu et al. SpWA: An efficient sparse winograd convolutional neural networks accelerator on FPGAs
US11442785B2 (en) Computation method and product thereof
CN111967468B (zh) 一种基于fpga的轻量级目标检测神经网络的实现方法
CN107301456B (zh) 基于向量处理器的深度神经网络多核加速实现方法
CN111062472B (zh) 一种基于结构化剪枝的稀疏神经网络加速器及其加速方法
Zhou et al. Transpim: A memory-based acceleration via software-hardware co-design for transformer
CN111626414B (zh) 一种动态多精度神经网络加速单元
Hong et al. Dfx: A low-latency multi-fpga appliance for accelerating transformer-based text generation
US20230026006A1 (en) Convolution computation engine, artificial intelligence chip, and data processing method
CN113469350B (zh) 一种适于npu的深度卷积神经网络加速方法和***
Asgari et al. Meissa: Multiplying matrices efficiently in a scalable systolic architecture
Que et al. A reconfigurable multithreaded accelerator for recurrent neural networks
Liu et al. Leveraging fine-grained structured sparsity for cnn inference on systolic array architectures
Li et al. Optimized data reuse via reordering for sparse matrix-vector multiplication on fpgas
Zhu et al. Taming unstructured sparsity on GPUs via latency-aware optimization
Kalali et al. Near-precise parameter approximation for multiple multiplications on a single dsp block
CN113240570B (zh) 一种GEMM运算加速器及基于GoogLeNet的图像处理加速方法
CN116167304B (zh) 基于神威架构的油藏数值模拟gmres优化方法及***
US11886347B2 (en) Large-scale data processing computer architecture
CN115222028A (zh) 基于fpga的一维cnn-lstm加速平台及实现方法
Shyamala et al. Design and implementation of GPU-based matrix chain multiplication using C++ AMP

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: 20230915

Address after: 518000, Zone 2111, Area A, 2nd Floor, Building R2-B, Gaoxin Industrial Village, No. 020 Gaoxin South Seventh Road, Gaoxin Community, Yuehai Street, Nanshan District, Shenzhen City, Guangdong Province

Patentee after: Shenzhen Xiangruilai Technology Co.,Ltd.

Address before: 510640 No. five, 381 mountain road, Guangzhou, Guangdong, Tianhe District

Patentee before: SOUTH CHINA University OF TECHNOLOGY

TR01 Transfer of patent right
TR01 Transfer of patent right

Effective date of registration: 20231027

Address after: B508, Unit 1, Building 6, Shenzhen Software Park, No. 2 Gaoxin Middle Road, Maling Community, Yuehai Street, Nanshan District, Shenzhen City, Guangdong Province, 518000

Patentee after: Shenzhen Aitesi Information Technology Co.,Ltd.

Address before: 518000, Zone 2111, Area A, 2nd Floor, Building R2-B, Gaoxin Industrial Village, No. 020 Gaoxin South Seventh Road, Gaoxin Community, Yuehai Street, Nanshan District, Shenzhen City, Guangdong Province

Patentee before: Shenzhen Xiangruilai Technology Co.,Ltd.