|
|
EDA365欢迎您登录！
您需要 登录 才可以下载或查看，没有帐号？注册
x
转—— DE1上跑完整的卷积神经网络
; z( K( K: e; a* ?0 S; K距离上次发帖有点久了，这段时间都在调试另一块Stratix V的板子。
* g' ]3 R0 B6 {! f. _" a5 F
2 u: ~! @- T) ?' a9 n7 F! Z闲话少絮，这里展示下项目最新进展。) v$ I- \9 e3 d) q2 v7 L6 V
; N3 G' H; F% Q7 R- M" E
前面的帖子（【Altera SoC体验之旅】+ 正式开启OpenCL模式）简要介绍了什么是神经网络，这次为大家详细展示一个利用卷积神经网络实现图片自动分类的例程。
$ A0 l' g4 S- y2 V9 h$ c* E& d
* ?' i, d: X" g神经网络的优点：自动从数据中学习经验知识，无需复杂的模型和算法。
( W+ N$ `4 b4 W+ x! m8 m. }$ a
3 X0 U# J: L U. }& A x9 L6 x缺点：有监督学习，需要大量的带标签数据；参数量太少时容易过拟合，泛化能力差，参数量太大时训练收敛很慢（有可能需要几个月到几年）。
2 ~ T& T4 o+ g0 Z4 j( _6 Z; B% r: b. e1 f8 }* A( o
为了克服上述缺点，人们发掘了各种计算资源，包括多核CPU、GPU、DSP、ASIC、FPGA，甚至使用模拟电路。: q) y0 m' V/ f6 Z: h I/ }& |* S+ V" b
, Z# z5 c1 M3 U# v+ P6 ?) A
使用CPU实现卷积神经网络比较方便调试，但性能太差，一般人们都选用更快的GPU实现。目前开源的框架大多都支持GPU，如伯克利大学Caffe（caffe官网）和Google Convnet（convnet2）。& A$ A3 R6 Q$ K' r# L" T
- P9 A0 H/ S+ Q+ @5 w+ ^$ ^
微软在今年2月宣布使用Stratix V完成了CNN加速器，处理 CIFAR10 图片速度可达每秒2300多张，链接如下（微软FPGA-CNN加速器）（技术报告）。0 V6 Y/ \8 @$ U" q+ k! C
$ t' v0 R9 p, _7 o' i G* C; K, r* y
这里我们也使用CIFAR10图片数据，在Cyclone V板子上跑一个卷积神经网络CNN demo。由于板子上计算资源太少（DSP Slice只有80多个），实现完整的网络不太现实，只能在FPGA上实现基本计算单元，然后由HPS统一调度。性能预期不会太高，后面给出。" G0 ?8 h. G( P5 E
% t# G7 ~* u9 e' R9 J% Q0 pCIFAR10图片都是什么呢？先来张图！4 }$ b8 _( h3 Y
! X, \+ v* ]3 w' Y, I& E; d
![]() + s! B/ R. y {7 N
2 c5 j' S. f1 ?, o有兴趣的朋友可以到官网下载（CIFAR10官网）。上面提到过，CNN是有监督学习系统，需要大量带label的数据，CIFAR10就是这样一个开放的数据库，提供了60000张不同类别的图片，分为10个类（如上图左侧所示），每个类别有600张图。这个数据集不算特别大，适合在嵌入式平台上实现。而更大的数据集有ImageNet-1000（ImageNet官网），拥有120多万张高清无码大图，我下载到硬盘，占用了近200GB空间（只能忍痛将其他rmvb和avi删掉了）！. Q4 x" ?7 u4 C9 \' n! N
9 e, T4 Z5 n/ q5 |
有朋友会问，不用这些数据行不行，我们的智能手机里面照片能不能用于CNN做训练？
3 h9 H O8 F8 e3 Y6 @6 f. s. H# q, E" ~) K3 X
答案是可以的，只是你的数据集很不“均匀”，采样不够“完备”，训练出的模型是真实模型的“有偏估计”，而上述两个数据集经过了种种考验，已经是学术界公认的优质数据集，一年一度的ILSVRC比赛就采用了这些数据集。
% ]* P# Y" M. }: V& W$ J3 b2 `
9 u8 {1 q9 H# ^7 | \6 D% [, Q x1 `2 m7 C; ?2 o# Z3 x5 D- J+ Z$ `
说完数据，再说模型。先来看一张经典的CNN结构：1 J) q0 ^/ L% C x7 A& n
![]() ( V4 v6 C y0 b9 g( S2 r
+ t/ n r, X/ S: ]这是世界上第一个将CNN实用化的例子，实现了手写体字母自动识别。在这个CNN模型中，可以看到输入是一张32 x 32的二维图像，经过卷积层(Convolution)、下采样层（Subsampling，也称Pooling)、全连接层（Full Connection，也称Inner Product)后，得到一组概率密度，我们选其中概率最大的元素作为该模型对输入图像的分类结果。所以实现CNN时，只需要实现三种基本算法：卷积、下采样、矩阵乘。除此之外，每层输出都可选择是否经过非线性变换，常用的非线性变换有ReLU和Sigmoid，前者计算较为简单，使用较为广泛。
8 h$ o6 o" H+ E& c8 {9 W8 i5 }* q
, ?& M/ n3 i) D# e% X% c# `. o0 d- o
7 ^/ F. M: C: i
Caffe框架中提供了专门为cifar10数据定制的模型，是用proto格式写的，我们的demo也基于这个模型。内容如下：
" ], |5 f$ P# d6 T. J# h# k# Z! A
6 F0 a$ l; o+ q* Q: _- name: "CIFAR10_quick_test"
- input: "data"
- input_dim: 1
- input_dim: 3
- input_dim: 32
- input_dim: 32
- layers {
- name: "conv1"
- type: CONVOLUTION
- bottom: "data"
- top: "conv1"
- blobs_lr: 1
- blobs_lr: 2
- convolution_param {
- num_output: 32
- pad: 2
- kernel_size: 5
- stride: 1
- }
- }
- layers {
- name: "pool1"
- type: POOLING
- bottom: "conv1"
- top: "pool1"
- pooling_param {
- pool: MAX
- kernel_size: 3
- stride: 2
- }
- }
- layers {
- name: "relu1"
- type: RELU
- bottom: "pool1"
- top: "pool1"
- }
- layers {
- name: "conv2"
- type: CONVOLUTION
- bottom: "pool1"
- top: "conv2"
- blobs_lr: 1
- blobs_lr: 2
- convolution_param {
- num_output: 32
- pad: 2
- kernel_size: 5
- stride: 1
- }
- }
- layers {
- name: "relu2"
- type: RELU
- bottom: "conv2"
- top: "conv2"
- }
- layers {
- name: "pool2"
- type: POOLING
- bottom: "conv2"
- top: "pool2"
- pooling_param {
- pool: AVE
- kernel_size: 3
- stride: 2
- }
- }
- layers {
- name: "conv3"
- type: CONVOLUTION
- bottom: "pool2"
- top: "conv3"
- blobs_lr: 1
- blobs_lr: 2
- convolution_param {
- num_output: 64
- pad: 2
- kernel_size: 5
- stride: 1
- }
- }
- layers {
- name: "relu3"
- type: RELU
- bottom: "conv3"
- top: "conv3"
- }
- layers {
- name: "pool3"
- type: POOLING
- bottom: "conv3"
- top: "pool3"
- pooling_param {
- pool: AVE
- kernel_size: 3
- stride: 2
- }
- }
- layers {
- name: "ip1"
- type: INNER_PRODUCT
- bottom: "pool3"
- top: "ip1"
- blobs_lr: 1
- blobs_lr: 2
- inner_product_param {
- num_output: 64
- }
- }
- layers {
- name: "ip2"
- type: INNER_PRODUCT
- bottom: "ip1"
- top: "ip2"
- blobs_lr: 1
- blobs_lr: 2
- inner_product_param {
- num_output: 10
- }
- }
- layers {
- name: "prob"
- type: SOFTMAX
- bottom: "ip2"
- top: "prob"
- }
+ ?9 ~: k. |# }$ \9 Y
+ M4 E! @$ f, Z
% a: l } P# D7 B" f+ @. g# s M7 X7 L v6 q& G0 a7 p4 a% ?' y d
8 p1 x+ v, W& f8 t
0 ~2 y; A/ U( v6 I& s* ]; i, P
可见，上述模型经过了3个卷积层（conv1, conv2, conv3），每个卷积层后面都跟着下采样层（pool1, pool2, pool3），之后有两个全连接层(ip1, ip2)，最后一层prob为SOFTMAX分类层，是计算概率密度的，这里我们不需要关心。
6 A3 [! _ E' N+ x0 V# g) J0 G& J/ |6 q0 q
下面三张图分别统计了CNN模型各层的参数量、数据量和计算量。# `* P/ R0 M1 G! B8 ~+ I2 h
! a3 P8 o/ F' M/ C
![]() ![]() ![]()
% G) m9 [5 ?/ ?, @* u: S6 `$ f/ ]3 D
可以看出，卷积层的参数量很少，但数据量很大；全连接层刚好相反，参数量较大，但数据量很少。
/ m- d6 ?; R) Q! A1 I$ [9 t9 P# K' T7 ^5 ^' o" m
通过计算量统计发现conv2计算量最大，其次是conv3和conv1。全连接层的计算量相对卷积层较小，但不可忽略。其他层（pool1, pool2以及各级relu）由于计算量太小，本设计中没有将其实现为Open CL kernel，而是直接CPU端实现。1 O1 A% K7 n" i7 @ X
3 z) {/ j/ X1 N( L! G综上所述，我们重点实现两个算法：卷积和矩阵乘，分别对应卷积层、全连接层的实现。
9 N1 W% o0 a' }- Z; S7 p" T! ?) Q, D$ R1 f! k" u' V& Q; |3 j4 \
在DE1-SOC上我利用了友晶提供的Open CL BSP，支持C语言开发FPGA。6 C1 E% f+ s" F
) E# q+ Z" U9 j4 z( r3 x: ?6 s卷积层计算kernel函数如下：
. _ C& ~. a. b0 J6 ?
$ b! A: R7 a! V3 B9 O0 t- __attribute__((num_compute_units(4)))
- __kernel
- void conv(__global float * a, __global float * b, __global float * c, const int M, const int N, const int K)
- {
- int gx = get_global_id(0);
- int gy = get_global_id(1);
- float tmp=0.0f;
- for(int x = 0; x < K; x ++)
- {
- for(int y = 0; y < K; y ++)
- {
- tmp += a[(gx + x) * M + (gy + y)] * b[x * K + y];
- }
- }
U: E" [; _( J6 z
6 t! F1 Z6 ~; \7 Z
9 C# f. Z$ I9 A3 Z! k- n
0 J2 N8 t( |( P7 k
7 ~$ I. w0 ] N" h7 L3 O, f# e3 Z8 N
全连接层计算采用矩阵乘实现，kernel函数如下：
4 J: S& `7 U1 \4 s. V7 [- d8 n. ^! E: R( J
- __attribute__((num_compute_units(4)))
- __kernel
- void gemm(__global float * a, __global float * b, __global float * c, const int M, const int N, const int K)
- {
- int gx = get_global_id(0);
- int gy = get_global_id(1);
- int sy = get_global_size(1);
- int sx = get_global_size(0);
- int s = sx * sy;
- for(int x = gx; x < M; x += sx)
- {
- for(int y = gy; y < N; y += sy)
- {
- float tmp=0.0f;
- for(int z = 0; z < K; z++)
- {
- tmp += a[z * M + x] * b[y * K + z];
- }
- c[y * M + x] = tmp;
- }
- }
- }
& Z2 g3 U; `. i; V; W$ U* I% D& T5 t2 u8 ?
1 J; t, n, K" P5 h& v
( R' ?( r8 d* T4 L+ ~5 T7 Z编译kernel函数需要使用Altera SDK for OpenCL，我用的版本是14.0.0.200，申请了两个月的license。编译使用命令行aoc，得到*.aocx文件。1 n9 H' A2 q4 L. F* o' K
, ~+ A3 [: Q2 b {' ]- ]Open CL编译输出报告中给出了资源占用情况：" x6 M$ r" j f$ h2 I: r% {
( w$ [' \! @1 u; f4 m% j! O
- +--------------------------------------------------------------------+
- ; Estimated Resource Usage Summary ;
- +----------------------------------------+---------------------------+
- ; Resource + Usage ;
- +----------------------------------------+---------------------------+
- ; Logic utilization ; 83% ;
- ; Dedicated logic registers ; 46% ;
- ; Memory blocks ; 57% ;
- ; DSP blocks ; 25% ;
- +----------------------------------------+---------------------------;* h4 I3 j# F$ m
- ~* {" A4 q$ d# R, {
! o$ h. L; J |6 r9 f可见，逻辑资源、存储器资源消耗较为明显，而DSP资源并未用尽，说明还有优化的空间。
/ M& k. x: G* B) q: i% M4 y1 r! j+ j" F$ }
编译主程序需要使用SoCEDS，我用的版本为14.0.2.274，也是命令行方式，在工程目录下执行make，结束后得到可执行文件cnn。( `9 k$ o1 P6 P$ @( x
: T2 n% _* K. {
将这两个文件拷贝到SD卡，按照前面的博客对板子进行设置，将CNN的模型、CIFAR10数据也拷贝到SD卡中，板子上电，mount SD卡到/mnt，执行cnn，得到输出如下：1 R0 o& Z, w0 t& s- [; w) M
) u6 d$ M: s4 r- ]2 B
- <div class="blockcode"><blockquote>Please input the number of images(1~100):100
- Loading data...OK!
- Constructing CNN...OK!
- Begin calculation...Elapsed Time = 141.861 s.
- Real Label = 3(cat), Calc Label = 3(cat), error count = 0
- Real Label = 8(ship), Calc Label = 8(ship), error count = 0
- Real Label = 8(ship), Calc Label = 8(ship), error count = 0
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 0
- Real Label = 6(frog), Calc Label = 6(frog), error count = 0
- Real Label = 6(frog), Calc Label = 6(frog), error count = 0
- Real Label = 1(automobile), Calc Label = 1(automobile), error count = 0
- Real Label = 6(frog), Calc Label = 6(frog), error count = 0
- Real Label = 3(cat), Calc Label = 3(cat), error count = 0
- Real Label = 1(automobile), Calc Label = 1(automobile), error count = 0
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 0
- Real Label = 9(truck), Calc Label = 9(truck), error count = 0
- Real Label = 5(dog), Calc Label = 5(dog), error count = 0
- Real Label = 7(horse), Calc Label = 7(horse), error count = 0
- Real Label = 9(truck), Calc Label = 9(truck), error count = 0
- Real Label = 8(ship), Calc Label = 8(ship), error count = 0
- Real Label = 5(dog), Calc Label = 5(dog), error count = 0
- Real Label = 7(horse), Calc Label = 7(horse), error count = 0
- Real Label = 8(ship), Calc Label = 8(ship), error count = 0
- Real Label = 6(frog), Calc Label = 6(frog), error count = 0
- Real Label = 7(horse), Calc Label = 7(horse), error count = 0
- Real Label = 0(airplane), Calc Label = 2(bird), error count = 1
- Real Label = 4(deer), Calc Label = 4(deer), error count = 1
- Real Label = 9(truck), Calc Label = 9(truck), error count = 1
- Real Label = 5(dog), Calc Label = 4(deer), error count = 2
- Real Label = 2(bird), Calc Label = 3(cat), error count = 3
- Real Label = 4(deer), Calc Label = 4(deer), error count = 3
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 3
- Real Label = 9(truck), Calc Label = 9(truck), error count = 3
- Real Label = 6(frog), Calc Label = 6(frog), error count = 3
- Real Label = 6(frog), Calc Label = 6(frog), error count = 3
- Real Label = 5(dog), Calc Label = 5(dog), error count = 3
- Real Label = 4(deer), Calc Label = 4(deer), error count = 3
- Real Label = 5(dog), Calc Label = 5(dog), error count = 3
- Real Label = 9(truck), Calc Label = 9(truck), error count = 3
- Real Label = 2(bird), Calc Label = 3(cat), error count = 4
- Real Label = 4(deer), Calc Label = 7(horse), error count = 5
- Real Label = 1(automobile), Calc Label = 9(truck), error count = 6
- Real Label = 9(truck), Calc Label = 9(truck), error count = 6
- Real Label = 5(dog), Calc Label = 5(dog), error count = 6
- Real Label = 4(deer), Calc Label = 4(deer), error count = 6
- Real Label = 6(frog), Calc Label = 6(frog), error count = 6
- Real Label = 5(dog), Calc Label = 5(dog), error count = 6
- Real Label = 6(frog), Calc Label = 6(frog), error count = 6
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 6
- Real Label = 9(truck), Calc Label = 9(truck), error count = 6
- Real Label = 3(cat), Calc Label = 5(dog), error count = 7
- Real Label = 9(truck), Calc Label = 9(truck), error count = 7
- Real Label = 7(horse), Calc Label = 7(horse), error count = 7
- Real Label = 6(frog), Calc Label = 6(frog), error count = 7
- Real Label = 9(truck), Calc Label = 9(truck), error count = 7
- Real Label = 8(ship), Calc Label = 8(ship), error count = 7
- Real Label = 0(airplane), Calc Label = 2(bird), error count = 8
- Real Label = 3(cat), Calc Label = 3(cat), error count = 8
- Real Label = 8(ship), Calc Label = 8(ship), error count = 8
- Real Label = 8(ship), Calc Label = 8(ship), error count = 8
- Real Label = 7(horse), Calc Label = 7(horse), error count = 8
- Real Label = 7(horse), Calc Label = 7(horse), error count = 8
- Real Label = 4(deer), Calc Label = 3(cat), error count = 9
- Real Label = 6(frog), Calc Label = 3(cat), error count = 10
- Real Label = 7(horse), Calc Label = 7(horse), error count = 10
- Real Label = 3(cat), Calc Label = 5(dog), error count = 11
- Real Label = 6(frog), Calc Label = 6(frog), error count = 11
- Real Label = 3(cat), Calc Label = 3(cat), error count = 11
- Real Label = 6(frog), Calc Label = 6(frog), error count = 11
- Real Label = 2(bird), Calc Label = 2(bird), error count = 11
- Real Label = 1(automobile), Calc Label = 1(automobile), error count = 11
- Real Label = 2(bird), Calc Label = 2(bird), error count = 11
- Real Label = 3(cat), Calc Label = 3(cat), error count = 11
- Real Label = 7(horse), Calc Label = 9(truck), error count = 12
- Real Label = 2(bird), Calc Label = 2(bird), error count = 12
- Real Label = 6(frog), Calc Label = 6(frog), error count = 12
- Real Label = 8(ship), Calc Label = 8(ship), error count = 12
- Real Label = 8(ship), Calc Label = 8(ship), error count = 12
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 12
- Real Label = 2(bird), Calc Label = 2(bird), error count = 12
- Real Label = 9(truck), Calc Label = 0(airplane), error count = 13
- Real Label = 3(cat), Calc Label = 3(cat), error count = 13
- Real Label = 3(cat), Calc Label = 2(bird), error count = 14
- Real Label = 8(ship), Calc Label = 8(ship), error count = 14
- Real Label = 8(ship), Calc Label = 8(ship), error count = 14
- Real Label = 1(automobile), Calc Label = 1(automobile), error count = 14
- Real Label = 1(automobile), Calc Label = 1(automobile), error count = 14
- Real Label = 7(horse), Calc Label = 7(horse), error count = 14
- Real Label = 2(bird), Calc Label = 2(bird), error count = 14
- Real Label = 5(dog), Calc Label = 7(horse), error count = 15
- Real Label = 2(bird), Calc Label = 2(bird), error count = 15
- Real Label = 7(horse), Calc Label = 7(horse), error count = 15
- Real Label = 8(ship), Calc Label = 8(ship), error count = 15
- Real Label = 9(truck), Calc Label = 9(truck), error count = 15
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 15
- Real Label = 3(cat), Calc Label = 4(deer), error count = 16
- Real Label = 8(ship), Calc Label = 8(ship), error count = 16
- Real Label = 6(frog), Calc Label = 6(frog), error count = 16
- Real Label = 4(deer), Calc Label = 4(deer), error count = 16
- Real Label = 6(frog), Calc Label = 6(frog), error count = 16
- Real Label = 6(frog), Calc Label = 6(frog), error count = 16
- Real Label = 0(airplane), Calc Label = 2(bird), error count = 17
- Real Label = 0(airplane), Calc Label = 0(airplane), error count = 17
- Real Label = 7(horse), Calc Label = 7(horse), error count = 17
- Classify Score = 83 %.
" c. ~( d& m& W4 D
- } n5 R/ `, D) k# D
2 X0 r* O! w7 `5 y! D( ~9 ]
' ~6 c1 h' y3 ^+ L2 P" i) b3 B9 A3 J2 V* ?% d" s7 N7 c/ r9 u
, S- c6 H @6 q" W$ ^9 t+ Z1 I/ q
上面的执行流程是这样的，首先输入测试样本数目（1到100），由于DE1板子FPGA端SDRAM容量较小，难以加载全部测试数据（10000张图片），故每次最多装入100张图片。之后载入数据到HPS内存，然后开始构建CNN模型，构建过程中也实现了Open CL的初始化。构建完毕，将输入图像依次通过CNN，得到一系列分类结果，与标签进行对比，统计错误分类个数，计算分类准确率。
6 S* T$ o7 I3 J, l+ e
8 f, q2 a* n* K$ k# `" P. Q经过测试，分类准确率达到83%，与Caffe测试结果一致。
3 @, c- H& p& P: U* T7 D- \, w/ r% v3 Q: `
6 k6 w, M! O% \# [8 ~
' k( D/ |1 D1 A0 _. I( T经过以上测试，可以得到结论：
5 b1 m$ @+ n. M% [$ i' n9 h) v) w. h1 c. {5 U5 H! b& k. I
（1）使用Open CL可以很方便地移植高级语言编写的算法；
/ J* m" Q1 O+ ?" u0 n) P
- @; I6 o% w# B- ^/ C' v% m（2）CNN在移植过程中需要考虑实际硬件，定制合适的模型和数据； h& |4 }( @( l( I/ S: o
! x/ h' c9 d# _（3）Cyclone 5逻辑资源较少（85K，Open CL kernel占用了83%），如果希望进一步提高计算速度，一方面可以选用高性能器件（如Stratix V、Arria 10），另一方面可以使用RTL自己搭建计算系统。8 C; p: R' b8 u2 t8 Q' p5 U' _
4 C! D$ V+ h1 s" e& |. E4 @
8 S# n# O; [" e+ p我目前正在Stratix V上用RTL搭建并行+流水线CNN加速器，有兴趣的童鞋可以进一步交流。% q. J O7 @% ~6 ^; D5 `! v# A% s
|
|
/1 
发表于 2019-4-15 07:00
收藏
淘帖
支持！
反对！