【Altera SoC体验之旅】+ DE1上跑完整的卷积神经网络

zhaoyongke

【Altera SoC体验之旅】+ DE1上跑完整的卷积神经网络 [复制链接]

 

距离上次发帖有点久了，这段时间都在调试另一块Stratix V的板子。

闲话少絮，这里展示下项目最新进展。

前面的帖子（【Altera SoC体验之旅】+ 正式开启OpenCL模式）简要介绍了什么是神经网络，这次为大家详细展示一个利用卷积神经网络实现图片自动分类的例程。

神经网络的优点：自动从数据中学习经验知识，无需复杂的模型和算法。

缺点：有监督学习，需要大量的带标签数据；参数量太少时容易过拟合，泛化能力差，参数量太大时训练收敛很慢（有可能需要几个月到几年）。

为了克服上述缺点，人们发掘了各种计算资源，包括多核CPU、GPU、DSP、ASIC、FPGA，甚至使用模拟电路。

使用CPU实现卷积神经网络比较方便调试，但性能太差，一般人们都选用更快的GPU实现。目前开源的框架大多都支持GPU，如伯克利大学Caffe（caffe官网）和Google Convnet（convnet2）。

微软在今年2月宣布使用Stratix V完成了CNN加速器，处理 CIFAR10 图片速度可达每秒2300多张，链接如下（微软FPGA-CNN加速器）（技术报告）。

这里我们也使用CIFAR10图片数据，在Cyclone V板子上跑一个卷积神经网络CNN demo。由于板子上计算资源太少（DSP Slice只有80多个），实现完整的网络不太现实，只能在FPGA上实现基本计算单元，然后由HPS统一调度。性能预期不会太高，后面给出。

CIFAR10图片都是什么呢？先来张图！



有兴趣的朋友可以到官网下载（CIFAR10官网）。上面提到过，CNN是有监督学习系统，需要大量带label的数据，CIFAR10就是这样一个开放的数据库，提供了60000张不同类别的图片，分为10个类（如上图左侧所示），每个类别有600张图。这个数据集不算特别大，适合在嵌入式平台上实现。而更大的数据集有ImageNet-1000（ImageNet官网），拥有120多万张高清无码大图，我下载到硬盘，占用了近200GB空间（只能忍痛将其他rmvb和avi删掉了）！

有朋友会问，不用这些数据行不行，我们的智能手机里面照片能不能用于CNN做训练？

答案是可以的，只是你的数据集很不“均匀”，采样不够“完备”，训练出的模型是真实模型的“有偏估计”，而上述两个数据集经过了种种考验，已经是学术界公认的优质数据集，一年一度的ILSVRC比赛就采用了这些数据集。


说完数据，再说模型。先来看一张经典的CNN结构：


这是世界上第一个将CNN实用化的例子，实现了手写体字母自动识别。在这个CNN模型中，可以看到输入是一张32 x 32的二维图像，经过卷积层(Convolution)、下采样层（Subsampling，也称Pooling)、全连接层（Full Connection，也称Inner Product)后，得到一组概率密度，我们选其中概率最大的元素作为该模型对输入图像的分类结果。所以实现CNN时，只需要实现三种基本算法：卷积、下采样、矩阵乘。除此之外，每层输出都可选择是否经过非线性变换，常用的非线性变换有ReLU和Sigmoid，前者计算较为简单，使用较为广泛。



Caffe框架中提供了专门为cifar10数据定制的模型，是用proto格式写的，我们的demo也基于这个模型。内容如下：

1. name: "CIFAR10_quick_test"
   
2. input: "data"
   
3. input_dim: 1
   
4. input_dim: 3
   
5. input_dim: 32
   
6. input_dim: 32
   
7. layers {
   
8.   name: "conv1"
   
9.   type: CONVOLUTION
   
10.   bottom: "data"
    
11.   top: "conv1"
    
12.   blobs_lr: 1
    
13.   blobs_lr: 2
    
14.   convolution_param {
    
15.     num_output: 32
    
16.     pad: 2
    
17.     kernel_size: 5
    
18.     stride: 1
    
19.   }
    
20. }
    
21. layers {
    
22.   name: "pool1"
    
23.   type: POOLING
    
24.   bottom: "conv1"
    
25.   top: "pool1"
    
26.   pooling_param {
    
27.     pool: MAX
    
28.     kernel_size: 3
    
29.     stride: 2
    
30.   }
    
31. }
    
32. layers {
    
33.   name: "relu1"
    
34.   type: RELU
    
35.   bottom: "pool1"
    
36.   top: "pool1"
    
37. }
    
38. layers {
    
39.   name: "conv2"
    
40.   type: CONVOLUTION
    
41.   bottom: "pool1"
    
42.   top: "conv2"
    
43.   blobs_lr: 1
    
44.   blobs_lr: 2
    
45.   convolution_param {
    
46.     num_output: 32
    
47.     pad: 2
    
48.     kernel_size: 5
    
49.     stride: 1
    
50.   }
    
51. }
    
52. layers {
    
53.   name: "relu2"
    
54.   type: RELU
    
55.   bottom: "conv2"
    
56.   top: "conv2"
    
57. }
    
58. layers {
    
59.   name: "pool2"
    
60.   type: POOLING
    
61.   bottom: "conv2"
    
62.   top: "pool2"
    
63.   pooling_param {
    
64.     pool: AVE
    
65.     kernel_size: 3
    
66.     stride: 2
    
67.   }
    
68. }
    
69. layers {
    
70.   name: "conv3"
    
71.   type: CONVOLUTION
    
72.   bottom: "pool2"
    
73.   top: "conv3"
    
74.   blobs_lr: 1
    
75.   blobs_lr: 2
    
76.   convolution_param {
    
77.     num_output: 64
    
78.     pad: 2
    
79.     kernel_size: 5
    
80.     stride: 1
    
81.   }
    
82. }
    
83. layers {
    
84.   name: "relu3"
    
85.   type: RELU
    
86.   bottom: "conv3"
    
87.   top: "conv3"
    
88. }
    
89. layers {
    
90.   name: "pool3"
    
91.   type: POOLING
    
92.   bottom: "conv3"
    
93.   top: "pool3"
    
94.   pooling_param {
    
95.     pool: AVE
    
96.     kernel_size: 3
    
97.     stride: 2
    
98.   }
    
99. }
    
100. layers {
     
101.   name: "ip1"
     
102.   type: INNER_PRODUCT
     
103.   bottom: "pool3"
     
104.   top: "ip1"
     
105.   blobs_lr: 1
     
106.   blobs_lr: 2
     
107.   inner_product_param {
     
108.     num_output: 64
     
109.   }
     
110. }
     
111. layers {
     
112.   name: "ip2"
     
113.   type: INNER_PRODUCT
     
114.   bottom: "ip1"
     
115.   top: "ip2"
     
116.   blobs_lr: 1
     
117.   blobs_lr: 2
     
118.   inner_product_param {
     
119.     num_output: 10
     
120.   }
     
121. }
     
122. layers {
     
123.   name: "prob"
     
124.   type: SOFTMAX
     
125.   bottom: "ip2"
     
126.   top: "prob"
     
127. }
     

复制代码

可见，上述模型经过了3个卷积层（conv1, conv2, conv3），每个卷积层后面都跟着下采样层（pool1, pool2, pool3），之后有两个全连接层(ip1, ip2)，最后一层prob为SOFTMAX分类层，是计算概率密度的，这里我们不需要关心。

下面三张图分别统计了CNN模型各层的参数量、数据量和计算量。



可以看出，卷积层的参数量很少，但数据量很大；全连接层刚好相反，参数量较大，但数据量很少。

通过计算量统计发现conv2计算量最大，其次是conv3和conv1。全连接层的计算量相对卷积层较小，但不可忽略。其他层（pool1, pool2以及各级relu）由于计算量太小，本设计中没有将其实现为Open CL kernel，而是直接CPU端实现。

综上所述，我们重点实现两个算法：卷积和矩阵乘，分别对应卷积层、全连接层的实现。

在DE1-SOC上我利用了友晶提供的Open CL BSP，支持C语言开发FPGA。

卷积层计算kernel函数如下：

1. __attribute__((num_compute_units(4)))
   
2. __kernel
   
3. void conv(__global float * a, __global float * b, __global float * c, const int M, const int N, const int K)
   
4. {
   
5.         int gx = get_global_id(0);
   
6.         int gy = get_global_id(1);
   
7. float tmp=0.0f;
   
8. for(int x = 0; x < K; x ++)
   
9.         {
   
10.                 for(int y = 0; y < K; y ++)
    
11.                 {
    
12. tmp += a[(gx + x) * M + (gy + y)] * b[x * K + y];
    
13.                         }
    
14.                 }
    

复制代码

全连接层计算采用矩阵乘实现，kernel函数如下：

1. __attribute__((num_compute_units(4)))
   
2. __kernel
   
3. void gemm(__global float * a, __global float * b, __global float * c, const int M, const int N, const int K)
   
4. {
   
5.         int gx = get_global_id(0);
   
6.         int gy = get_global_id(1);
   
7.         int sy = get_global_size(1);
   
8.         int sx = get_global_size(0);
   
9. 
   
10.         int s = sx * sy;
    
11.         for(int x = gx; x < M; x += sx)
    
12.         {
    
13.                 for(int y = gy; y < N; y += sy)
    
14.                 {
    
15.                         float tmp=0.0f;
    
16.                         for(int z = 0; z < K; z++)
    
17.                         {
    
18.                                 tmp += a[z * M + x] * b[y * K + z];
    
19.                         }
    
20.                         c[y * M + x] = tmp;
    
21.                 }
    
22.         }
    
23. }
    

复制代码

编译kernel函数需要使用Altera SDK for OpenCL，我用的版本是14.0.0.200，申请了两个月的license。编译使用命令行aoc，得到*.aocx文件。

Open CL编译输出报告中给出了资源占用情况：

1. +--------------------------------------------------------------------+
   
2. ; Estimated Resource Usage Summary                                   ;
   
3. +----------------------------------------+---------------------------+
   
4. ; Resource                               + Usage                     ;
   
5. +----------------------------------------+---------------------------+
   
6. ; Logic utilization                      ;   83%                     ;
   
7. ; Dedicated logic registers              ;   46%                     ;
   
8. ; Memory blocks                          ;   57%                     ;
   
9. ; DSP blocks                             ;   25%                     ;
   
10. +----------------------------------------+---------------------------;

复制代码

可见，逻辑资源、存储器资源消耗较为明显，而DSP资源并未用尽，说明还有优化的空间。

编译主程序需要使用SoCEDS，我用的版本为14.0.2.274，也是命令行方式，在工程目录下执行make，结束后得到可执行文件cnn。

将这两个文件拷贝到SD卡，按照前面的博客对板子进行设置，将CNN的模型、CIFAR10数据也拷贝到SD卡中，板子上电，mount SD卡到/mnt，执行cnn，得到输出如下：

1. <div class="blockcode"><blockquote>Please input the number of images(1~100):100
   
2. Loading data...OK!
   
3. Constructing CNN...OK!
   
4. Begin calculation...Elapsed Time = 141.861 s.
   
5. Real Label = 3(cat), Calc Label = 3(cat), error count = 0
   
6. Real Label = 8(ship), Calc Label = 8(ship), error count = 0
   
7. Real Label = 8(ship), Calc Label = 8(ship), error count = 0
   
8. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 0
   
9. Real Label = 6(frog), Calc Label = 6(frog), error count = 0
   
10. Real Label = 6(frog), Calc Label = 6(frog), error count = 0
    
11. Real Label = 1(automobile), Calc Label = 1(automobile), error count = 0
    
12. Real Label = 6(frog), Calc Label = 6(frog), error count = 0
    
13. Real Label = 3(cat), Calc Label = 3(cat), error count = 0
    
14. Real Label = 1(automobile), Calc Label = 1(automobile), error count = 0
    
15. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 0
    
16. Real Label = 9(truck), Calc Label = 9(truck), error count = 0
    
17. Real Label = 5(dog), Calc Label = 5(dog), error count = 0
    
18. Real Label = 7(horse), Calc Label = 7(horse), error count = 0
    
19. Real Label = 9(truck), Calc Label = 9(truck), error count = 0
    
20. Real Label = 8(ship), Calc Label = 8(ship), error count = 0
    
21. Real Label = 5(dog), Calc Label = 5(dog), error count = 0
    
22. Real Label = 7(horse), Calc Label = 7(horse), error count = 0
    
23. Real Label = 8(ship), Calc Label = 8(ship), error count = 0
    
24. Real Label = 6(frog), Calc Label = 6(frog), error count = 0
    
25. Real Label = 7(horse), Calc Label = 7(horse), error count = 0
    
26. Real Label = 0(airplane), Calc Label = 2(bird), error count = 1
    
27. Real Label = 4(deer), Calc Label = 4(deer), error count = 1
    
28. Real Label = 9(truck), Calc Label = 9(truck), error count = 1
    
29. Real Label = 5(dog), Calc Label = 4(deer), error count = 2
    
30. Real Label = 2(bird), Calc Label = 3(cat), error count = 3
    
31. Real Label = 4(deer), Calc Label = 4(deer), error count = 3
    
32. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 3
    
33. Real Label = 9(truck), Calc Label = 9(truck), error count = 3
    
34. Real Label = 6(frog), Calc Label = 6(frog), error count = 3
    
35. Real Label = 6(frog), Calc Label = 6(frog), error count = 3
    
36. Real Label = 5(dog), Calc Label = 5(dog), error count = 3
    
37. Real Label = 4(deer), Calc Label = 4(deer), error count = 3
    
38. Real Label = 5(dog), Calc Label = 5(dog), error count = 3
    
39. Real Label = 9(truck), Calc Label = 9(truck), error count = 3
    
40. Real Label = 2(bird), Calc Label = 3(cat), error count = 4
    
41. Real Label = 4(deer), Calc Label = 7(horse), error count = 5
    
42. Real Label = 1(automobile), Calc Label = 9(truck), error count = 6
    
43. Real Label = 9(truck), Calc Label = 9(truck), error count = 6
    
44. Real Label = 5(dog), Calc Label = 5(dog), error count = 6
    
45. Real Label = 4(deer), Calc Label = 4(deer), error count = 6
    
46. Real Label = 6(frog), Calc Label = 6(frog), error count = 6
    
47. Real Label = 5(dog), Calc Label = 5(dog), error count = 6
    
48. Real Label = 6(frog), Calc Label = 6(frog), error count = 6
    
49. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 6
    
50. Real Label = 9(truck), Calc Label = 9(truck), error count = 6
    
51. Real Label = 3(cat), Calc Label = 5(dog), error count = 7
    
52. Real Label = 9(truck), Calc Label = 9(truck), error count = 7
    
53. Real Label = 7(horse), Calc Label = 7(horse), error count = 7
    
54. Real Label = 6(frog), Calc Label = 6(frog), error count = 7
    
55. Real Label = 9(truck), Calc Label = 9(truck), error count = 7
    
56. Real Label = 8(ship), Calc Label = 8(ship), error count = 7
    
57. Real Label = 0(airplane), Calc Label = 2(bird), error count = 8
    
58. Real Label = 3(cat), Calc Label = 3(cat), error count = 8
    
59. Real Label = 8(ship), Calc Label = 8(ship), error count = 8
    
60. Real Label = 8(ship), Calc Label = 8(ship), error count = 8
    
61. Real Label = 7(horse), Calc Label = 7(horse), error count = 8
    
62. Real Label = 7(horse), Calc Label = 7(horse), error count = 8
    
63. Real Label = 4(deer), Calc Label = 3(cat), error count = 9
    
64. Real Label = 6(frog), Calc Label = 3(cat), error count = 10
    
65. Real Label = 7(horse), Calc Label = 7(horse), error count = 10
    
66. Real Label = 3(cat), Calc Label = 5(dog), error count = 11
    
67. Real Label = 6(frog), Calc Label = 6(frog), error count = 11
    
68. Real Label = 3(cat), Calc Label = 3(cat), error count = 11
    
69. Real Label = 6(frog), Calc Label = 6(frog), error count = 11
    
70. Real Label = 2(bird), Calc Label = 2(bird), error count = 11
    
71. Real Label = 1(automobile), Calc Label = 1(automobile), error count = 11
    
72. Real Label = 2(bird), Calc Label = 2(bird), error count = 11
    
73. Real Label = 3(cat), Calc Label = 3(cat), error count = 11
    
74. Real Label = 7(horse), Calc Label = 9(truck), error count = 12
    
75. Real Label = 2(bird), Calc Label = 2(bird), error count = 12
    
76. Real Label = 6(frog), Calc Label = 6(frog), error count = 12
    
77. Real Label = 8(ship), Calc Label = 8(ship), error count = 12
    
78. Real Label = 8(ship), Calc Label = 8(ship), error count = 12
    
79. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 12
    
80. Real Label = 2(bird), Calc Label = 2(bird), error count = 12
    
81. Real Label = 9(truck), Calc Label = 0(airplane), error count = 13
    
82. Real Label = 3(cat), Calc Label = 3(cat), error count = 13
    
83. Real Label = 3(cat), Calc Label = 2(bird), error count = 14
    
84. Real Label = 8(ship), Calc Label = 8(ship), error count = 14
    
85. Real Label = 8(ship), Calc Label = 8(ship), error count = 14
    
86. Real Label = 1(automobile), Calc Label = 1(automobile), error count = 14
    
87. Real Label = 1(automobile), Calc Label = 1(automobile), error count = 14
    
88. Real Label = 7(horse), Calc Label = 7(horse), error count = 14
    
89. Real Label = 2(bird), Calc Label = 2(bird), error count = 14
    
90. Real Label = 5(dog), Calc Label = 7(horse), error count = 15
    
91. Real Label = 2(bird), Calc Label = 2(bird), error count = 15
    
92. Real Label = 7(horse), Calc Label = 7(horse), error count = 15
    
93. Real Label = 8(ship), Calc Label = 8(ship), error count = 15
    
94. Real Label = 9(truck), Calc Label = 9(truck), error count = 15
    
95. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 15
    
96. Real Label = 3(cat), Calc Label = 4(deer), error count = 16
    
97. Real Label = 8(ship), Calc Label = 8(ship), error count = 16
    
98. Real Label = 6(frog), Calc Label = 6(frog), error count = 16
    
99. Real Label = 4(deer), Calc Label = 4(deer), error count = 16
    
100. Real Label = 6(frog), Calc Label = 6(frog), error count = 16
     
101. Real Label = 6(frog), Calc Label = 6(frog), error count = 16
     
102. Real Label = 0(airplane), Calc Label = 2(bird), error count = 17
     
103. Real Label = 0(airplane), Calc Label = 0(airplane), error count = 17
     
104. Real Label = 7(horse), Calc Label = 7(horse), error count = 17
     
105. Classify Score = 83 %.
     

复制代码

上面的执行流程是这样的，首先输入测试样本数目（1到100），由于DE1板子FPGA端SDRAM容量较小，难以加载全部测试数据（10000张图片），故每次最多装入100张图片。之后载入数据到HPS内存，然后开始构建CNN模型，构建过程中也实现了Open CL的初始化。构建完毕，将输入图像依次通过CNN，得到一系列分类结果，与标签进行对比，统计错误分类个数，计算分类准确率。

经过测试，分类准确率达到83%，与Caffe测试结果一致。



经过以上测试，可以得到结论：

（1）使用Open CL可以很方便地移植高级语言编写的算法；

（2）CNN在移植过程中需要考虑实际硬件，定制合适的模型和数据；

（3）Cyclone 5逻辑资源较少（85K，Open CL kernel占用了83%），如果希望进一步提高计算速度，一方面可以选用高性能器件（如Stratix V、Arria 10），另一方面可以使用RTL自己搭建计算系统。


我目前正在Stratix V上用RTL搭建并行+流水线CNN加速器，有兴趣的童鞋可以进一步交流。





码字不易，申请加精。

Mr_Hertz

学习了，支持一下！

zhaoyongke

Mr_Hertz 发表于 2015-5-30 19:47
学习了，支持一下！

感谢支持！

574433742

好厉害的样子，我要学习一下，，然后在模仿一下。。。

zhaoyongke

574433742 发表于 2015-5-31 10:48
好厉害的样子，我要学习一下，，然后在模仿一下。。。

共同进步

zgbkdlm

略吊~

zhaoyongke

zgbkdlm 发表于 2015-6-1 14:53
略吊~

chenzhufly

不明觉厉！前来学习

zhaoyongke

chenzhufly 发表于 2015-6-1 17:24
不明觉厉！前来学习

期待柱版详细讲解HPS-FPGA高速AXI通信

chenzhufly

这个alex不是开坛授课了吗
我也在努力的学习中呢

zhaoyongke

chenzhufly 发表于 2015-6-1 21:26
这个alex不是开坛授课了吗
我也在努力的学习中呢

原来已经发了，还没看

qq365317289

楼主好厉害~方便的话可以留个邮箱交流一下把~

ZT_15

      博主您好！我是一名大四学生，专业是微电子，我的毕业设计是将CNN算法在FPGA上实现加速。前段时间忙于考研复试、调剂的事情，半个月以前才开始做毕设。通过查找各种资料，我对CNN算法思想已经掌握，并用Matlab语言实现了CNN算法，可以进行简单的图片识别。我了解到CNN算法在卷积层和全连接层运算量非常大，也是CNN算法在FPGA上实现加速需要重点考虑的地方，但是我是第一次接触OpenCL语言，将CNN卷积层和全连接层用OpenCL语言实现对我有很大难度。事关我能否毕业，我现在真的非常着急。今天我意外发现这个论坛，看到博主最近也在研究这个问题，感觉真是拨云见雾。博主好人，能不能将CNN算法的OpenCL全部实现代码发给我，我保证尊重博主的科研成果，不会发给别人。我知道我这样做非常不好，但是，我真的是没有办法了。我是新注册用户，不能跟博主私聊，如果博主看到了我的回复，请发发慈悲。我的邮箱是782097656@qq.com 。我会非常感激博主！

zhaoyongke

qq365317289 发表于 2015-6-5 16:12
楼主好厉害~方便的话可以留个邮箱交流一下把~

已发私信

zhaoyongke

ZT_15 发表于 2015-6-5 16:18
博主您好！我是一名大四学生，专业是微电子，我的毕业设计是将CNN算法在FPGA上实现加速。前段时间忙于考研复试、调剂的事情，半个月以前才开始做毕设。通过查找各种资料，我对CNN算法思想已经掌握，并用Matlab语言实现了CNN算法，可以进行简单的图片识别。我了解到CNN算法在卷积层和全连接层运算量非常大，也是CNN算法在FPGA上实现加速需要重点考虑的地方，但是我是第一次接触OpenCL语言，将CNN卷积层和全连接层用OpenCL语言实现对我有很大难度。事关我能否毕业，我现在真的非常着急。今天我意外发现这个论坛，看到博主最近也在研究这个问题，感觉真是拨云见雾。博主好人，能不能将CNN算法的OpenCL全部实现代码发给我，我保证尊重博主的科研成果，不会发给别人。我知道我这样做非常不好，但是，我真的是没有办法了。我是新注册用户，不能跟博主私聊，如果博主看到了我的回复，请发发慈悲。我的邮箱是 。我会非常感激博主！

不好意思，目前还不能公开

guaiguaidou

请问一下opencl怎么申请license？

zhaoyongke

guaiguaidou 发表于 2015-6-14 10:42
请问一下opencl怎么申请license？

找站内管理员申请即可

ZT_15

zhaoyongke 发表于 2015-6-10 14:40
不好意思，目前还不能公开

感谢博主百忙之中回复我

guaiguaidou

zhaoyongke 发表于 2015-6-15 09:03
找站内管理员申请即可

我的板是买的，也能申请吗

调戏、和尚/

                 
       哇塞。数据量好大。得好好消化一下才行。。。