CN111340224B - 适用于低资源嵌入式芯片的cnn网络的加速设计方法 - Google Patents
适用于低资源嵌入式芯片的cnn网络的加速设计方法 Download PDFInfo
- Publication number
- CN111340224B CN111340224B CN202010125198.1A CN202010125198A CN111340224B CN 111340224 B CN111340224 B CN 111340224B CN 202010125198 A CN202010125198 A CN 202010125198A CN 111340224 B CN111340224 B CN 111340224B
- Authority
- CN
- China
- Prior art keywords
- data
- calculation
- sram
- layer
- convolution
- 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
Links
- 238000000034 method Methods 0.000 title claims abstract description 33
- 238000013461 design Methods 0.000 title claims abstract description 12
- 238000004364 calculation method Methods 0.000 claims abstract description 78
- 238000011176 pooling Methods 0.000 claims abstract description 27
- 229910052754 neon Inorganic materials 0.000 claims abstract description 25
- GKAOGPIIYCISHV-UHFFFAOYSA-N neon atom Chemical compound [Ne] GKAOGPIIYCISHV-UHFFFAOYSA-N 0.000 claims abstract description 25
- 238000005457 optimization Methods 0.000 claims description 8
- 238000012546 transfer Methods 0.000 claims description 2
- 230000001133 acceleration Effects 0.000 abstract description 7
- 238000013528 artificial neural network Methods 0.000 abstract description 4
- 238000013135 deep learning Methods 0.000 abstract description 4
- 238000009825 accumulation Methods 0.000 description 6
- 238000012545 processing Methods 0.000 description 5
- 238000013527 convolutional neural network Methods 0.000 description 4
- 230000003139 buffering effect Effects 0.000 description 2
- 230000001186 cumulative effect Effects 0.000 description 2
- 238000010586 diagram Methods 0.000 description 2
- 239000000463 material Substances 0.000 description 2
- 230000009286 beneficial effect Effects 0.000 description 1
- 230000000295 complement effect Effects 0.000 description 1
- 239000012141 concentrate Substances 0.000 description 1
- 238000007796 conventional method Methods 0.000 description 1
- 230000000694 effects Effects 0.000 description 1
- 230000003993 interaction Effects 0.000 description 1
- 230000009191 jumping Effects 0.000 description 1
- 229910044991 metal oxide Inorganic materials 0.000 description 1
- 150000004706 metal oxides Chemical class 0.000 description 1
- 238000012986 modification Methods 0.000 description 1
- 230000004048 modification Effects 0.000 description 1
- 238000011160 research Methods 0.000 description 1
- 239000004065 semiconductor Substances 0.000 description 1
- 238000000926 separation method Methods 0.000 description 1
- 230000003068 static effect Effects 0.000 description 1
Classifications
-
- 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/08—Learning methods
- G06N3/082—Learning methods modifying the architecture, e.g. adding, deleting or silencing nodes or connections
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F13/00—Interconnection of, or transfer of information or other signals between, memories, input/output devices or central processing units
- G06F13/14—Handling requests for interconnection or transfer
- G06F13/20—Handling requests for interconnection or transfer for access to input/output bus
- G06F13/28—Handling requests for interconnection or transfer for access to input/output bus using burst mode transfer, e.g. direct memory access DMA, cycle steal
-
- 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
-
- Y—GENERAL TAGGING OF NEW TECHNOLOGICAL DEVELOPMENTS; GENERAL TAGGING OF CROSS-SECTIONAL TECHNOLOGIES SPANNING OVER SEVERAL SECTIONS OF THE IPC; TECHNICAL SUBJECTS COVERED BY FORMER USPC CROSS-REFERENCE ART COLLECTIONS [XRACs] AND DIGESTS
- Y02—TECHNOLOGIES OR APPLICATIONS FOR MITIGATION OR ADAPTATION AGAINST CLIMATE CHANGE
- Y02D—CLIMATE CHANGE MITIGATION TECHNOLOGIES IN INFORMATION AND COMMUNICATION TECHNOLOGIES [ICT], I.E. INFORMATION AND COMMUNICATION TECHNOLOGIES AIMING AT THE REDUCTION OF THEIR OWN ENERGY USE
- Y02D10/00—Energy efficient computing, e.g. low power processors, power management or thermal management
Landscapes
- Engineering & Computer Science (AREA)
- Theoretical Computer Science (AREA)
- Physics & Mathematics (AREA)
- General Physics & Mathematics (AREA)
- General Engineering & Computer Science (AREA)
- Data Mining & Analysis (AREA)
- Computing Systems (AREA)
- Computational Linguistics (AREA)
- Biomedical Technology (AREA)
- Evolutionary Computation (AREA)
- General Health & Medical Sciences (AREA)
- Molecular Biology (AREA)
- Biophysics (AREA)
- Artificial Intelligence (AREA)
- Life Sciences & Earth Sciences (AREA)
- Mathematical Physics (AREA)
- Software Systems (AREA)
- Health & Medical Sciences (AREA)
- Complex Calculations (AREA)
Abstract
本发明公开了一种适用于低资源嵌入式芯片的CNN网络的加速设计方法，涉及神经网络深度学习技术领域。本发明包括如下步骤：对卷积/池化运算的计算方向进行重新的设计，采用以通道方向优先的计算顺序；对CNN网络中卷积/池化层的空间分配和计算流程进行优化；利用Neon指令集来提升计算性能。本发明充分利用现有嵌入式芯片中所普遍存在的DMA、SRAM、Neon等资源，通过对CNN网络中卷积单元和池化单元的计算步骤进行优化编排和SRAM空间重复利用，将数据搬移和数值计算分离，用DMA来搬数据，用Neon来实现数值计算，中间结果不写DDR等手段，实现在低资源的嵌入式芯片中可以更快更高效的运行CNN网络的目的。
Description
技术领域
本发明属于神经网络深度学习技术领域，特别是涉及一种适用于低资源嵌入式芯片的CNN网络的加速设计方法。
背景技术
卷积神经网络(Convolutional Neural Network,CNN)是深度学习技术领域中的一项重要的创新。卷积神经网络作为一种典型的多层神经网络始终处于研究的核心地位。采用的局部连接和权值共享的方式，一方面减少了权值的数量使得网络易于优化，另一方面降低了模型的复杂度，也就是减小了过拟合的风险。因此，CNN被应用在很多的机器视觉相关的任务中,并取得了很大的成功。
CNN带来的优异的效果，激发了很多存量的嵌入式设备厂商通过软件升级的方式来对其原有产品进行智能赋能的需求。而当前主流的CNN的计算过程往往需要借助于一些专门的硬件资源才可以实现。比如，目前主要是借助于GPU、FPGA、ASIC这三种硬件资源才能够实现部署。那么，在计算资源和内存资源都比较贫乏的嵌入式芯片上(低资源)如何把CNN网络部署上去，让它尽可能快速高效的运行起来，哪怕是部署相对简化的CNN网络，也成为一个值得研究的问题。
目前，按照常规方法，如果要在嵌入式芯片运行CNN网络，主要是仅利用了现有的CPU资源和DDR资源。针对CNN网络中的每一层卷积运算，基本流程如下：CPU读取DDR中的输入数据，做卷积/池化计算，将中间结果回存到DDR中。因为我们知道即使简化版的CNN网络层数也是很多的。所以，如果一直上面的常规流程处理。那么对CPU来说，既被用来处理数据搬移又被用来处理数值计算。对DDR来说，既用来存放输入输出层数据又被用来存放中间层数据。显然，这是一种比较低效的实现方式。这种低效的实现方案也就产生两个明显的问题：
1、本就计算速度不快CPU还得不到充分利用；
2、本就不充足的DDR带宽资源，还增加了很多额外的数据读写。
本发明针对如何在低资源嵌入式芯片上运行CNN网络，提出了一种高效的优化设计方案。
发明内容
本发明的目的在于提供一种适用于低资源嵌入式芯片的CNN网络的加速设计方法，充分挖掘现有芯片的能力，避免背景技术提到的两个明显问题，使得CNN网络在低资源嵌入式芯片上部署运行成为可能。
为解决上述技术问题，本发明是通过以下技术方案实现的：
本发明为一种适用于低资源嵌入式芯片的CNN网络的加速设计方法，包括如下步骤：
S001：对卷积/池化运算的计算方向进行重新的设计，采用以通道方向优先的计算顺序；
S002：对CNN网络中卷积/池化层的空间分配和计算流程进行优化；
S003：利用Neon指令集来提升计算性能；
充分利用嵌入式芯片中普遍存在的资源(DMA、片内SRAM、Neon)，设计了一套高效的执行编排方案，采用以通道方向优先的计算顺序，以一行来作为最小的计算单位，将数据的搬移过程和计算过程分离，利用DMA实现串行计算过程并行化，利用SRAM作为缓冲减少DDR带宽消耗利用Neon指令对计算进行加速，实现在低资源的嵌入式芯片中运行一些简化版的CNN网络的目的，给传统的嵌入式设备提供了更多的智能解决方案。
进一步地，所述步骤S001中采用以通道方向优先的计算顺序，具体计算顺序为：通道→宽度→高度，以一行作为最小的计算单位；只要满足一行的卷积/池化计算的数据要求就可以进行计算，同时只需载入一行数据，使得利用片内SRAM作为数据缓冲成为可能。
进一步地，所述步骤S002中空间分配的优化：对片内SRAM地址空间进行统一的分配管理，以一行作为数据的缓冲单位，大大减少每次独立计算所需要的数据缓冲区；针对CNN中的中间层数据，行计算的结果直接缓冲于片内SRAM中，不需要输出到DDR。
进一步地，所述步骤S002中计算流程的优化：采用由后往前的执行顺序来驱动每一层的运算；除了输入层和输出层的数据需要搬移到DDR中外，网络中的其他中间层的数据都是临时的存储在SRAM中，每层最多只临时存放一个卷积行/池化行的数据量，用完即被滚动回收，使得整个SRAM的空间在计算过程中可以得到重复利用的。
进一步地，所述步骤S003中利用Neon指令集来提升计算性能的方式为：使用Neon指令集中的装载、移位、乘加、存储指令实现卷积运算和池化运算，在一次指令周期内同时处理8个数据点，使得性能相比于常规ARM指令提升了5倍以上。
本发明具有以下有益效果：
本发明充分利用现有嵌入式芯片中所普遍存在的DMA、SRAM、Neon等资源，通过对CNN网络中卷积单元和池化单元的计算步骤进行优化编排和SRAM空间重复利用，将数据搬移和数值计算分离，用DMA来搬数据，用Neon来实现数值计算，中间结果不写DDR等手段，实现在低资源的嵌入式芯片中可以更快更高效的运行CNN网络的目的。
当然，实施本发明的任一产品并不一定需要同时达到以上所述的所有优点。
附图说明
为了更清楚地说明本发明实施例的技术方案，下面将对实施例描述所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本发明的一些实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据这些附图获得其他的附图。
图1为本发明的计算优先顺序的示意图；
图2为本发明的数据分配和数据搬移的流程图；
图3为本发明的计算驱动流程图；
图4为本发明的Neon加速的卷积处理示意图。
具体实施方式
下面将结合本发明实施例中的附图，对本发明实施例中的技术方案进行清楚、完整地描述，显然，所描述的实施例仅仅是本发明一部分实施例，而不是全部的实施例。基于本发明中的实施例，本领域普通技术人员在没有作出创造性劳动前提下所获得的所有其它实施例，都属于本发明保护的范围。
请参阅图1-4所示，本发明为一种适用于低资源嵌入式芯片的CNN网络的加速设计方法，包括如下步骤：
S001：对卷积/池化运算的计算方向进行重新的设计，采用以通道方向优先的计算顺序，具体计算顺序为：通道→宽度→高度，以一行作为最小的计算单位；只要满足一行的卷积/池化计算的数据要求就可以进行计算，同时只需载入一行数据，使得利用片内SRAM作为数据缓冲成为可能。常规的卷积和池化运算是按照：宽度→高度→通道这样的计算顺序进行的，这样的顺序一次卷积/池化计算要求载入一整块数据，所需要的内存消耗比较大；而采用以通道方向优先的计算顺序(通道→宽度→高度)，以一行来作为最小的计算单位，进行卷积/池化计算，这样就不需要等到所有的输入数据全部搬移到位，只要满足一行的卷积/池化计算要求，就可以对该行进行计算，这样的设计，一次仅需要载入一行数据，使得利用片内SRAM作为数据缓冲成为可能，也使得后续的数据搬移过程和数值计算过程分离成为可能。
S002：对CNN网络中卷积/池化层的空间分配和计算流程进行优化：对片内SRAM地址空间进行统一的分配管理，以一行作为数据的缓冲单位，大大减少每次独立计算所需要的数据缓冲区；空间分配的优化：针对CNN中的中间层数据，行计算的结果直接缓冲于片内SRAM中，而不需要输出到DDR。首先，对片内SRAM地址空间进行统一的分配管理，对每层之间的计算关联关系进行解耦，对运算过程进行并行化处理，降低对DRR的数据的吞吐需求，如步骤S001中：采用以通道方向优先的计算顺序，以一行作为最小的计算单位进行卷积/池化计算，这样也就意味着以一行来作为数据的缓冲单位，使得每次独立的计算单位所需要的数据缓冲区的大小大大节省。针对CNN中的中间层数据，行计算的结果可以直接缓冲于片内SRAM中，而不需要输出到DDR。同时因为片内SRAM带宽资源充足，还可以节省因额外的数据搬移所产生的对DDR带宽的消耗。
其次，计算流程的优化：采用由后往前的执行顺序来驱动每一层的运算；除了输入层和输出层的数据需要搬移到DDR中外，网络中的其他中间层的数据都是临时的存储在SRAM中，每层最多只临时存放一个卷积行/池化行的数据量，用完即被滚动回收，使得整个SRAM的空间在计算过程中可以得到重复利用的，根据步骤S002中对片内SRAM地址空间的统一分配管理。
再次，数据的搬移过程和计算过程分离，利用DMA来做数据搬移，利用CPU来专注于做数值计算，节省了CPU等待数据的时间。实现DMA和CPU同时工作，做到并行化加速的优化目的。利用DMA来实现数据搬移，主要包括两个部分：1、将输入数据从DDR按行搬移输入到分配的SRAM地址空间中；2、将计算结果从片内SRAM按行搬移输出到DDR中。这样CPU就可以不用管数据的搬移,而专心于数值计算,就实现了并行化处理，提升了CPU的计算利用率。
S003：利用Neon指令集来提升计算性能。Neon指令集是ARM系统中的一种SIMD(Single Instruction Multiple Data)架构，一次指令可以处理多个数据，使用这些Neon指令可以实现卷积运算和池化运算，可以提高CPU的计算性能。在本发明中，利用Neon指令集来提升计算性能的方式为：使用Neon指令集中的装载、移位、乘加、存储等指令实现卷积运算和池化运算，在一次指令周期内同时处理8个数据点，使得性能相比于常规ARM指令提升了5倍以上。
充分利用嵌入式芯片中普遍存在的资源(DMA、片内SRAM、Neon)，设计了一套高效的执行编排方案，采用以通道方向优先的计算顺序，以一行来作为最小的计算单位，将数据的搬移过程和计算过程分离，利用DMA实现串行计算过程并行化，利用SRAM作为缓冲减少DDR带宽消耗利用Neon指令对计算进行加速，实现在低资源的嵌入式芯片中运行一些简化版的CNN网络的目的，给传统的嵌入式设备提供了更多的智能解决方案。
具体实施方式一：
S1：读取神经网络结构文件，分析每一个层的数据大小和卷积核大小，计算确认每一个层所需要的数据缓冲空间，然后在SRAM中为其分配对应的地址范围。
以深度学习计算框架Caffe为例，分析对应的prototxt文件。根据网络结构，获取每一层的数据大小的形态信息(记作：C，H，W)以及每一层的卷积核的大小(记作：Kh，Kw)。其中，第1层需要预留第Kh+1片用来让DMA提前预搬数据，第2层到第n层的数据只需要有K片之后就可以进行运算并回收了，那么具体每一个层的内存消耗标记如下：
输入层：通过DMA把数据从DDR搬移到SRAM中，每次搬移的DMA操作单元为(W*C)，需要SRAM的大小记为：Sram_Size＝(Kh+1)*(W*C)；
中间层：当上一级数据准备好Kh片后就可以进行运算，需要SRAM的大小记为：Sram_Size＝Kh*(W*C)；
输出层：通过DMA把片内SRAM的数据送到外部DDR中去，需要SRAM的大小记为：Sram_Size＝2*(W*C)，即两DMA操作单元以乒乓缓冲的方式进行搬移；
在确定每一层的计算单元所需要的内存大小后，为其在片内SRAM中分配对应的地址空间。具体的数据分布情况如下图2中所示，其中，深色部分为SRAM内部为数值计算所用的数据区，其中，浅色部分为DMA数据搬移所用的和DDR交互的数据区。
S2：采用由后往前的执行顺序来驱动每一层的数值计算和数据搬移，实现CNN并行化加速，如图3中所示，处理策略描述如下：
输出层：如果自己的缓冲列表中有效的缓冲块(计算已完成)，则启动DMA数据搬移，将有效数据搬移输出到DRR中；
中间层：如果其前一层的缓冲列表(对当前层来说就是数据输入)已经有Kh个有效数据行，则启动CPU进行数值计算；否则，结束当次循环，跳到输入层；
输入层：如果自己的缓冲列表中有空闲的缓冲块，则启动DMA搬移，将数据从DDR中搬移输入到SRAM，直到没有空闲块为止；
程序按照上述处理流程重复执行驱动这个CNN网络的计算。
S3：用Neon指令来对卷积操作和池化操作进行加速。
具体处理流程如图4中所示，其中前缀为i的表示输入数据，前缀为k的表示卷积核，前缀为o的表示输出数据。使用neon指令依次读取8个数据点和8个卷积核到寄存器中，然后进行乘加运算，再将8个点的计算结果输出到缓冲区中。摘录优化指令如下：
1.#装载参数
2.int8x8_t k11＝vld1_dup_s8(k11)
3.int8x8_t k12＝vld1_dup_s8(k12)
4.int8x8_t k13＝vld1_dup_s8(k13)
5.int8x8_t k21＝vld1_dup_s8(k21)
6.int8x8_t k22＝vld1_dup_s8(k22)
7.int8x8_t k23＝vld1_dup_s8(k23)
8.int8x8_t k31＝vld1_dup_s8(k31)
9.int8x8_t k32＝vld1_dup_s8(k32)
10.int8x8_t k33＝vld1_dup_s8(k33)
11.LAB:一次处理多个像素点
12.#临时结果缓冲区
13.int16x8_t sum＝vld1q_s16(sum_cache_line)
14.#因为乘法指令要两个周期
15.int8x8_t val_a＝vld1_s8(i11)
16.int16x8_t mul_a＝vmull_s8(val_a,k11)
17.int8x8_t val_b＝vld1_s8(i12)
18.int16x8_t mul_b＝vmull_s8(val_b,k12)
19.#右移常数位再累积(右移3bit)
20.mul_a＝vshrq_n_s16(mul_a,3)
21.sum＝vaddq_s16(sum,mul_a)#累积11
22.val_a＝vld1_s8(i13)
23.mul_a＝vmull_s8(val_a,k13)
24.mul_b＝vshrq_n_s16(mul_b,3)
25.sum＝vqaddq_s16(sum,mul_b)#累积12
26.val_b＝vld1_s8(i21)
27.mul_b＝vmull_s8(val_b,k21)
28.mul_a＝vshrq_n_s16(mul_a,3)
29.sum＝vaddq_s16(sum,mul_a)#累积13
30.val_a＝vld1_s8(i22)
31.mul_a＝vmull_s8(val_a,k22)
32.mul_b＝vshrq_n_s16(mul_b,3)
33.sum＝vqaddq_s16(sum,mul_b)#累积21
34.val_b＝vld1_s8(i23)
35.mul_b＝vmull_s8(val_b,k23)
36.mul_a＝vshrq_n_s16(mul_a,3)
37.sum＝vaddq_s16(sum,mul_a)#累积22
38.val_a＝vld1_s8(i31)
39.mul_a＝vmull_s8(val_a,k31)
40.mul_b＝vshrq_n_s16(mul_b,3)
41.sum＝vqaddq_s16(sum,mul_b)#累积23
42.val_b＝vld1_s8(i32)
43.mul_b＝vmull_s8(val_b,k32)
44.mul_a＝vshrq_n_s16(mul_a,3)
45.sum＝vaddq_s16(sum,mul_a)#累积31
46.val_a＝vld1_s8(i33)
47.mul_a＝vmull_s8(val_a,k33)
48.mul_b＝vshrq_n_s16(mul_b,3)
49.sum＝vqaddq_s16(sum,mul_b)#累积32
50.mul_a＝vshrq_n_s16(mul_a,3)
51.sum＝vaddq_s16(sum,mul_a)#累积33
52.#结果写入缓冲区中
53.vst1q_s16(sum_cache_line,sum)；
54.处理下组8个点:
本发明充分利用嵌入式芯片中普遍存在的资源(DMA、片内SRAM、Neon)，设计了一套高效的执行编排方案，采用以通道方向优先的计算顺序，以一行来作为最小的计算单位，将数据的搬移过程和计算过程分离，利用DMA实现串行计算过程并行化，利用SRAM作为缓冲减少DDR带宽消耗利用Neon指令对计算进行加速，实现在低资源的嵌入式芯片中运行一些简化版的CNN网络的目的，给传统的嵌入式设备提供了更多的智能解决方案。
在本说明书的描述中，参考术语“一个实施例”、“示例”、“具体示例”等的描述意指结合该实施例或示例描述的具体特征、结构、材料或者特点包含于本发明的至少一个实施例或示例中。在本说明书中，对上述术语的示意性表述不一定指的是相同的实施例或示例。而且，描述的具体特征、结构、材料或者特点可以在任何的一个或多个实施例或示例中以合适的方式结合。
以上公开的本发明优选实施例只是用于帮助阐述本发明。优选实施例并没有详尽叙述所有的细节，也不限制该发明仅为所述的具体实施方式。显然，根据本说明书的内容，可作很多的修改和变化。本说明书选取并具体描述这些实施例，是为了更好地解释本发明的原理和实际应用，从而使所属技术领域技术人员能很好地理解和利用本发明。本发明仅受权利要求书及其全部范围和等效物的限制。
Claims (2)
1.适用于低资源嵌入式芯片的CNN网络的加速设计方法，其特征在于，包括如下步骤：
S001：对卷积/池化运算的计算方向进行重新的设计，采用以通道方向优先的计算顺序；
S002：对CNN网络中卷积/池化层的空间分配和计算流程进行优化；
S003：利用Neon指令集来提升计算性能；
所述步骤S001中采用以通道方向优先的计算顺序，具体计算顺序为：通道→宽度→高度，以一行作为最小的计算单位；只要满足一行的卷积/池化计算的数据要求就可以进行计算，同时只需载入一行数据；
所述步骤S002中空间分配的优化：对片内SRAM地址空间进行统一的分配管理，以一行作为数据的缓冲单位，减少每次独立计算所需要的数据缓冲区；针对CNN中的中间层数据，行计算的结果直接缓冲于片内SRAM中，不需要输出到DDR；根据网络结构，获取每一层的数据大小的形态信息，记作：C，H，W，以及每一层的卷积核的大小，记作：Kh，Kw；其中，第1层需要预留第Kh+1片用来让DMA提前预搬数据，第2层到第n层的数据只需要有K片之后就可以进行运算并回收，每一个层的内存消耗标记如下：
输入层：通过DMA把数据从DDR搬移到SRAM中，每次搬移的DMA操作单元为W*C，需要SRAM的大小记为：Sram_Size＝(Kh+1)*(W*C)；
中间层：当上一级数据准备好Kh片后就可以进行运算，需要SRAM的大小记为：Sram_Size＝Kh*(W*C)；
输出层：通过DMA把片内SRAM的数据送到外部DDR中去，需要SRAM的大小记为：Sram_Size＝2*(W*C)；
在确定每一层的计算单元所需要的内存大小后，为其在片内SRAM中分配对应的地址空间；
所述步骤S002中计算流程的优化：采用由后往前的执行顺序来驱动每一层的运算；除了输入层和输出层的数据需要搬移到DDR中外，网络中的其他中间层的数据都是临时的存储在SRAM中，每层最多只临时存放一个卷积行/池化行的数据量，用完即被滚动回收。
2.根据权利要求1所述的适用于低资源嵌入式芯片的CNN网络的加速设计方法，其特征在于，所述步骤S003中利用Neon指令集来提升计算性能的方式为：使用Neon指令集中的装载、移位、乘加、存储指令实现卷积运算和池化运算，在一次指令周期内同时处理8个数据点。
Priority Applications (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202010125198.1A CN111340224B (zh) | 2020-02-27 | 2020-02-27 | 适用于低资源嵌入式芯片的cnn网络的加速设计方法 |
Applications Claiming Priority (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202010125198.1A CN111340224B (zh) | 2020-02-27 | 2020-02-27 | 适用于低资源嵌入式芯片的cnn网络的加速设计方法 |
Publications (2)
| Publication Number | Publication Date |
|---|---|
| CN111340224A CN111340224A (zh) | 2020-06-26 |
| CN111340224B true CN111340224B (zh) | 2023-11-21 |
Family
ID=71187087
Family Applications (1)
| Application Number | Title | Priority Date | Filing Date |
|---|---|---|---|
| CN202010125198.1A Active CN111340224B (zh) | 2020-02-27 | 2020-02-27 | 适用于低资源嵌入式芯片的cnn网络的加速设计方法 |
Country Status (1)
| Country | Link |
|---|---|
| CN (1) | CN111340224B (zh) |
Families Citing this family (1)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| TWI835244B (zh) * | 2022-08-16 | 2024-03-11 | 聯陽半導體股份有限公司 | 計算裝置、計算裝置的操作方法及單晶片系統 |
Citations (11)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN108171317A (zh) * | 2017-11-27 | 2018-06-15 | 北京时代民芯科技有限公司 | 一种基于soc的数据复用卷积神经网络加速器 |
| EP3346427A1 (en) * | 2017-01-04 | 2018-07-11 | STMicroelectronics Srl | Configurable accelerator framework, system and method |
| CN108416434A (zh) * | 2018-02-07 | 2018-08-17 | 复旦大学 | 针对神经网络的卷积层与全连接层进行加速的电路结构 |
| CN108985450A (zh) * | 2018-06-28 | 2018-12-11 | 中国人民解放军国防科技大学 | 面向向量处理器的卷积神经网络运算向量化方法 |
| CN109615066A (zh) * | 2019-01-30 | 2019-04-12 | 新疆爱华盈通信息技术有限公司 | 一种针对neon优化的卷积神经网络的裁剪方法 |
| CN109934339A (zh) * | 2019-03-06 | 2019-06-25 | 东南大学 | 一种基于一维脉动阵列的通用卷积神经网络加速器 |
| CN110276444A (zh) * | 2019-06-04 | 2019-09-24 | 北京清微智能科技有限公司 | 基于卷积神经网络的图像处理方法及装置 |
| CN110399971A (zh) * | 2019-07-03 | 2019-11-01 | Oppo广东移动通信有限公司 | 一种卷积神经网络加速方法及装置、存储介质 |
| CN110705687A (zh) * | 2019-09-05 | 2020-01-17 | 北京三快在线科技有限公司 | 卷积神经网络硬件计算装置及方法 |
| CN110738308A (zh) * | 2019-09-23 | 2020-01-31 | 陈小柏 | 一种神经网络加速器 |
| CN110796235A (zh) * | 2019-10-21 | 2020-02-14 | 中国人民解放军国防科技大学 | 卷积神经网络Valid卷积的向量化实现方法 |
Family Cites Families (4)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| US10489703B2 (en) * | 2015-05-20 | 2019-11-26 | Nec Corporation | Memory efficiency for convolutional neural networks operating on graphics processing units |
| US10387740B2 (en) * | 2016-10-10 | 2019-08-20 | Gyrfalcon Technology Inc. | Object detection and recognition apparatus based on CNN based integrated circuits |
| AR109623A1 (es) * | 2018-02-16 | 2019-01-09 | Pescarmona Enrique Menotti | Proceso y sistema de análisis y gestión hidrológica para cuencas |
| US11487846B2 (en) * | 2018-05-04 | 2022-11-01 | Apple Inc. | Performing multiply and accumulate operations in neural network processor |
-
2020
- 2020-02-27 CN CN202010125198.1A patent/CN111340224B/zh active Active
Patent Citations (11)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| EP3346427A1 (en) * | 2017-01-04 | 2018-07-11 | STMicroelectronics Srl | Configurable accelerator framework, system and method |
| CN108171317A (zh) * | 2017-11-27 | 2018-06-15 | 北京时代民芯科技有限公司 | 一种基于soc的数据复用卷积神经网络加速器 |
| CN108416434A (zh) * | 2018-02-07 | 2018-08-17 | 复旦大学 | 针对神经网络的卷积层与全连接层进行加速的电路结构 |
| CN108985450A (zh) * | 2018-06-28 | 2018-12-11 | 中国人民解放军国防科技大学 | 面向向量处理器的卷积神经网络运算向量化方法 |
| CN109615066A (zh) * | 2019-01-30 | 2019-04-12 | 新疆爱华盈通信息技术有限公司 | 一种针对neon优化的卷积神经网络的裁剪方法 |
| CN109934339A (zh) * | 2019-03-06 | 2019-06-25 | 东南大学 | 一种基于一维脉动阵列的通用卷积神经网络加速器 |
| CN110276444A (zh) * | 2019-06-04 | 2019-09-24 | 北京清微智能科技有限公司 | 基于卷积神经网络的图像处理方法及装置 |
| CN110399971A (zh) * | 2019-07-03 | 2019-11-01 | Oppo广东移动通信有限公司 | 一种卷积神经网络加速方法及装置、存储介质 |
| CN110705687A (zh) * | 2019-09-05 | 2020-01-17 | 北京三快在线科技有限公司 | 卷积神经网络硬件计算装置及方法 |
| CN110738308A (zh) * | 2019-09-23 | 2020-01-31 | 陈小柏 | 一种神经网络加速器 |
| CN110796235A (zh) * | 2019-10-21 | 2020-02-14 | 中国人民解放军国防科技大学 | 卷积神经网络Valid卷积的向量化实现方法 |
Non-Patent Citations (2)
| Title |
|---|
| Tianqi Chen等.Training Deep Nets with Sublinear Memory Cost.《Machine Learning》.2016,1-10页. * |
| 吴焕 ; 吴俊敏 ; .基于Caffe加速卷积神经网络前向推理.计算机工程与设计.2018,(12),94-99. * |
Also Published As
| Publication number | Publication date |
|---|---|
| CN111340224A (zh) | 2020-06-26 |
Similar Documents
| Publication | Publication Date | Title |
|---|---|---|
| JP7329533B2 (ja) | 演算を加速するための方法および加速器装置 | |
| EP3557425B1 (en) | Accelerator and system for accelerating operations | |
| KR102572705B1 (ko) | 스케일러블 뉴럴 네트워크 프로세싱 엔진 | |
| CN112487750B (zh) | 一种基于存内计算的卷积加速计算系统及方法 | |
| WO2022206556A1 (zh) | 图像数据的矩阵运算方法、装置、设备及存储介质 | |
| US11163686B2 (en) | Method and apparatus for accessing tensor data | |
| WO2020073801A1 (zh) | 一种3d图像处理中数据读写方法及系统、存储介质及终端 | |
| WO2022078400A1 (zh) | 一种对多维数据进行处理的设备、方法和计算机程序产品 | |
| CN116107754A (zh) | 一种面向深度神经网络的内存管理方法及系统 | |
| CN118152980A (zh) | 一种分叉算子融合方法、装置、设备及存储介质 | |
| CN118132156B (zh) | 一种算子执行方法、设备、存储介质及程序产品 | |
| CN111340224B (zh) | 适用于低资源嵌入式芯片的cnn网络的加速设计方法 | |
| CN111475205B (zh) | 一种基于数据流解耦合的粗粒度可重构阵列结构设计方法 | |
| CN116431562B (zh) | 一种基于加速处理器的多头注意力机制融合计算分配方法 | |
| Shahbahrami et al. | FPGA implementation of parallel histogram computation | |
| CN109948787B (zh) | 用于神经网络卷积层的运算装置、芯片及方法 | |
| CN115878188A (zh) | 一种基于sve指令集的池化层函数的高性能实现方法 | |
| CN112433847B (zh) | 一种OpenCL内核提交的方法及装置 | |
| CN115222028A (zh) | 基于fpga的一维cnn-lstm加速平台及实现方法 | |
| CN118171710B (zh) | 一种稀疏矩阵乘法的npu加速方法 | |
| WO2023115529A1 (zh) | 芯片内的数据处理方法及芯片 | |
| US20160283235A1 (en) | Controlling data flow between processors in a processing system | |
| CN114820630B (zh) | 一种基于fpga的目标跟踪算法模型管道加速方法以及电路 | |
| CN114139107B (zh) | 池化装置与池化方法 | |
| US20220327391A1 (en) | Global pooling method for neural network, and many-core system |
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 | ||
| CB02 | Change of applicant information | ||
| CB02 | Change of applicant information |
Address after: 311400 4th floor, building 9, Yinhu innovation center, No.9 Fuxian Road, Yinhu street, Fuyang District, Hangzhou City, Zhejiang Province Applicant after: Zhejiang Xinmai Microelectronics Co.,Ltd. Address before: 311400 4th floor, building 9, Yinhu innovation center, No.9 Fuxian Road, Yinhu street, Fuyang District, Hangzhou City, Zhejiang Province Applicant before: Hangzhou xiongmai integrated circuit technology Co.,Ltd. |
|
| GR01 | Patent grant | ||
| GR01 | Patent grant |