海思3559AV100上运行自己训练的yolov3

最新推荐文章于 2022-09-13 17:14:14 发布
最新推荐文章于 2022-09-13 17:14:14 发布
阅读量2.6k 收藏 24
点赞数
分类专栏： 海思芯片
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/pursuit_zhangyu/article/details/107705995
版权

1、简介

现在海思芯片上只支持将caffemodel转成wk文件，然而网上也没有caffe-yolov3，只能使用darknet-yolov3进行训练，然后将模型文件转换成caffemodel文件，再使用ruyistdio转成指令集的wk文件，就可以在板子上运行了。期间遇到的坑做一下记录。

遇到了一个是段错误（Segmentation fault）和精度损失问题，下面提供了解决方法

2、遇到的问题（价值所在）

2.1 出现Segmentation fault

将自己训练的darknet模型转成wk文件，在海思芯片上运行，只修改了SAMPLE_SVP_NNIE_Yolov3_SoftwareInit里面的 pstSoftWareParam->u32ClassNum，会出现下面的错误

解决方法

还要修改sample_svp_nnie_software.h中关于yolov3的参数设置，修改为5+classnum（自己类别的个数）

2.2 yolov3出现精度损失问题

主要原因是生成的wk参数没有设置对，在ruyistdio中将image_type中的RGB_order设置为RGB，如下图

2.3 在仿真需要修改prototxt文件

input: "data"
input_dim: 1
input_dim: 3
input_dim: 416
input_dim: 416

修改为

input: "data"
input_shape {
	dim: 1
	dim: 3
	dim: 416
	dim: 416
}

3、Darknet训练自己的数据集

4、将darknet模型转成caffemodel

在darknet训练完以后需要将模型转变成caffemodel文件。这个过程需要安装 pycaffe环境，同时还需要给caffe环境添加不支持层 upsample。

4.1 pycaffe环境的安装

https://blog.csdn.net/pursuit_zhangyu/article/details/104770610

4.2 添加不支持层upsample重新编译

下面的文件我提供了下载

链接：https://pan.baidu.com/s/1z07aqFuUpkFYPHmEf9TWtA 
提取码：yykw

增加include/caffe/layers/upsample_layer.hpp

#ifndef CAFFE_UPSAMPLE_LAYER_HPP_
#define CAFFE_UPSAMPLE_LAYER_HPP_
 
#include <vector>
 
#include "caffe/blob.hpp"
#include "caffe/layer.hpp"
#include "caffe/proto/caffe.pb.h"
 
namespace caffe {
 
template <typename Dtype>
class UpsampleLayer : public Layer<Dtype> {
 public:
  explicit UpsampleLayer(const LayerParameter& param)
      : Layer<Dtype>(param) {}
  virtual void LayerSetUp(const vector<Blob<Dtype>*>& bottom,
      const vector<Blob<Dtype>*>& top);
  virtual void Reshape(const vector<Blob<Dtype>*>& bottom,
      const vector<Blob<Dtype>*>& top);
 
  virtual inline const char* type() const { return "Upsample"; }
  virtual inline int MinBottomBlobs() const { return 1; }
  virtual inline int MaxBottomBlobs() const { return 1; }
  virtual inline int ExactNumTopBlobs() const { return 1; }
 
  virtual void Forward_cpu(const vector<Blob<Dtype>*>& bottom,
      const vector<Blob<Dtype>*>& top);
  virtual void Forward_gpu(const vector<Blob<Dtype>*>& bottom,
      const vector<Blob<Dtype>*>& top);
  virtual void Backward_cpu(const vector<Blob<Dtype>*>& top,
      const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
  virtual void Backward_gpu(const vector<Blob<Dtype>*>& top,
      const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom);
 
 private:
  int scale_;
};
 
 
 
}  // namespace caffe
 
#endif  // CAFFE_UPSAMPLE_LAYER_HPP_

增加src/caffe/layers/upsample_layer.cpp

#include <vector>
#include "caffe/layers/upsample_layer.hpp"
 
namespace caffe {
 
template <typename Dtype>
void UpsampleLayer<Dtype>::LayerSetUp(
  const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
  UpsampleParameter upsample_param = this->layer_param_.upsample_param();
  scale_ = upsample_param.scale();
}
 
template <typename Dtype>
void UpsampleLayer<Dtype>::Reshape(
  const vector<Blob<Dtype>*>& bottom, const vector<Blob<Dtype>*>& top) {
  vector<int> out_shape;
  for (int i = 0; i < bottom[0]->num_axes(); i++) {
    out_shape.push_back(bottom[0]->shape(i));
  }
 
  out_shape[bottom[0]->num_axes() - 1] *= scale_;
  out_shape[bottom[0]->num_axes() - 2] *= scale_;
  top[0]->Reshape(out_shape);
}
 
template <typename Dtype>
void UpsampleLayer<Dtype>::Forward_cpu(const vector<Blob<Dtype>*>& bottom,
    const vector<Blob<Dtype>*>& top) {
 
  int N = top[0]->shape(0);
  int C = top[0]->shape(1);
  int H = top[0]->shape(2);
  int W = top[0]->shape(3);
 
  const Dtype *input = bottom[0]->cpu_data();
  Dtype *output = top[0]->mutable_cpu_data();
  for (int n = 0; n < N; n++) {
    for (int c = 0; c < C; c++) {
      for (int h = 0; h < H; h++) {
        for (int w = 0; w < W; w++) {
          int nw = w/scale_;
          int nh = h/scale_;
          int out_idx = (((n * C + c) * H) + h) * W + w;
          int in_idx = (((n * C + c) * (H / scale_)) + nh) * (W / scale_) + nw;
          output[out_idx] = input[in_idx];
        }
      }
    }
  }
}
 
template <typename Dtype>
void UpsampleLayer<Dtype>::Backward_cpu(const vector<Blob<Dtype>*>& top,
    const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
  int N = bottom[0]->shape(0);
  int C = bottom[0]->shape(1);
  int H = bottom[0]->shape(2);
  int W = bottom[0]->shape(3);
  const Dtype *output_grad = top[0]->cpu_diff();
  Dtype* bottom_diff = bottom[0]->mutable_cpu_diff();
  caffe_set(bottom[0]->count(), Dtype(0), bottom_diff);
  for (int n = 0; n < N; n++) {
    for (int c = 0; c < C; c++) {
      for (int h = 0; h < H; h++) {
        for (int w = 0; w < W; w++) {
          for (int i = 0; i < scale_; i++) {
            for (int j = 0; j < scale_; j++) {
              int nw = w * scale_ + i;
              int nh = h * scale_ + j;
              int out_idx = (((n * C + c) * H) + h) * W + w;
              int in_idx = (((n * C + c) * (H * scale_))
                  + nh) * (W * scale_) + nw;
              bottom_diff[out_idx] += output_grad[in_idx];
            }
          }
        }
      }
    }
  }
}
 
#ifdef CPU_ONLY
STUB_GPU(UpsampleLayer);
#endif
 
INSTANTIATE_CLASS(UpsampleLayer);
REGISTER_LAYER_CLASS(Upsample);
 
}  // namespace caffe

增加src/caffe/layers/upsample_layer.cu

#include <vector>
 
#include "caffe/filler.hpp"
#include "caffe/layers/upsample_layer.hpp"
#include "caffe/util/math_functions.hpp"
 
namespace caffe {
 
__device__ int translate_idx(int ii, int d1, int d2, int d3, int scale_factor) {
  int x, y, z, w;
  w = ii % d3;
  ii = ii/d3;
  z = ii % d2;
  ii = ii/d2;
  y = ii % d1;
  ii = ii/d1;
  x = ii;
  w = w/scale_factor;
  z = z/scale_factor;
  d2 /= scale_factor;
  d3 /= scale_factor;
  return (((x*d1+y)*d2)+z)*d3+w;
}
 
__device__ int translate_idx_inv(
    int ii, int d1, int d2, int d3, int scale_factor, int off_x, int off_y) {
  int x, y, z, w;
  w = ii % d3;
  ii = ii/d3;
  z = ii % d2;
  ii = ii/d2;
  y = ii % d1;
  ii = ii/d1;
  x = ii;
  w = w*scale_factor+off_x;
  z = z*scale_factor+off_y;
  d2 *= scale_factor;
  d3 *= scale_factor;
  return (((x*d1+y)*d2)+z)*d3+w;
}
 
template <typename Dtype>
__global__ void upscale(const Dtype *input, Dtype *output,
        int no_elements, int scale_factor, int d1, int d2, int d3) {
  int ii = threadIdx.x + blockDim.x * blockIdx.x;
  if (ii >= no_elements) return;
  int ipidx = translate_idx(ii, d1, d2, d3, scale_factor);
  output[ii]=input[ipidx];
}
 
template <typename Dtype>
__global__ void downscale(Dtype *gradInput_data, const Dtype *gradOutput_data,
                          int no_elements, int scale_factor, int d1, int d2,
                          int d3) {
  int ii = threadIdx.x + blockDim.x * blockIdx.x;
  if (ii >= no_elements) return;
  for (int i = 0; i < scale_factor; i++) {
    for (int j = 0; j < scale_factor; j++) {
      int ipidx = translate_idx_inv(ii, d1, d2, d3, scale_factor, i, j);
      gradInput_data[ii] += gradOutput_data[ipidx];
    }
  }
}
 
 
 
template <typename Dtype>
void UpsampleLayer<Dtype>::Forward_gpu(const vector<Blob<Dtype>*>& bottom,
      const vector<Blob<Dtype>*>& top) {
  int d1, d2, d3;
 
  d1 = top[0]->shape(1);
  d2 = top[0]->shape(2);
  d3 = top[0]->shape(3);
 
  int no_elements = top[0]->count();
 
  upscale<Dtype>  // NOLINT_NEXT_LINE(whitespace/operators)
      <<<CAFFE_GET_BLOCKS(no_elements), CAFFE_CUDA_NUM_THREADS>>>(
      bottom[0]->gpu_data(),
      top[0]->mutable_gpu_data(), no_elements, scale_, d1, d2, d3);
}
 
template <typename Dtype>
void UpsampleLayer<Dtype>::Backward_gpu(const vector<Blob<Dtype>*>& top,
      const vector<bool>& propagate_down, const vector<Blob<Dtype>*>& bottom) {
  int d1, d2, d3;
  Dtype* bottom_diff = bottom[0]->mutable_gpu_diff();
  d1 = bottom[0]->shape(1);
  d2 = bottom[0]->shape(2);
  d3 = bottom[0]->shape(3);
  int no_elements = bottom[0]->count();
  caffe_gpu_set(bottom[0]->count(), Dtype(0), bottom_diff);
  downscale<Dtype>  // NOLINT_NEXT_LINE(whitespace/operators)
      <<<CAFFE_GET_BLOCKS(no_elements), CAFFE_CUDA_NUM_THREADS>>>(
      bottom_diff, top[0]->gpu_diff(), no_elements, scale_, d1, d2, d3);
}
 
INSTANTIATE_LAYER_GPU_FUNCS(UpsampleLayer);
 
}  // namespace caffe

最后在src/caffe/proto/caffe.proto中增加

message LayerParameter {
 
           ……
 
  optional UpsampleParameter upsample_param = 148;
}
 
 
message UpsampleParameter {
  optional int32 scale = 1 [default = 1];
}

重新编译

make clean
make all -j
make py -j

4.3 转换工具

我将python2的代码改成了python3的了，下面我提供百度下载链接

链接：https://pan.baidu.com/s/1D133dWqyc71ZvXVGtOnKrA 
提取码：bzvb

python yolov3_2caffe.py darknet_model/yolov3_coco_416.cfg darknet_model/yolov3_coco_416.weights caffemodel/yolov3.prototxt caffemodel/yolov3.caffemodel

注意：我设置的是使用第2个gpu

4.4 可以测试转到caffe有没有精度损失

我测过没有任何损失问题

5、生成wk文件

 

6、仿真

如果需要仿真的话，如果是两个类别的话，作如下修改

仿真的步骤

1.首先build clean

2.release 编译

3.仿真

7、在板子上进行测试

代码修改，主要修改三个方面：

SAMPLE_SVP_NNIE_YOLOV3_EACH_BBOX_INFER_RESULT_NUM 是5+2(classnum）=7

如果是部分类别检测不出来sample_svp_nnie_software.c中找到static HI_S32 SVP_NNIE_Yolov3_GetResult，修改里面的

 

使用的SDK下载地址

链接：https://pan.baidu.com/s/1i3AjsZp_YfvyC2sQxRbSDQ 
提取码：ciei

pursuit_zhangyu
关注
专栏目录
华为海思 AI 芯片 (Hi3559A V100) 算法开发（四） 在 PC 仿真库使用 YOLOv3 进行图片目标检测以及 NMS、YOLO 讲解
随意聊点
05-06 5295
摘要 经过上一篇文章，大家编译运行 PC 仿真后，可能觉得有点迷惘，实际上你已经在后台完成了图像识别，保存结果，绘制识别框等一系列过程。这篇文章就教教大家怎么看识别的结果已经如何手动将识别框绘制出来。 仿真代码浅析 为了方便实际操作，这篇文章使用 VS 来快速运行，我觉得 vs 运行比 Ruyi 稍微快一点 为了速度再快一点，我们修改我们要运行的文件夹下...
海思NNIE开发（一）：海思Hi3559AV100/Hi3519AV100 NNIE深度学习模块开发与调试记录
热门推荐
夜风里唱的专栏
10-20 3万+
海思NNIE开发系列文章： 海思NNIE开发（一）：海思Hi3559AV100/Hi3519AV100 NNIE深度学习模块开发与调试记录 海思NNIE开发（二）：FasterRCNN在海思NNIE平台上的执行流程（一） 海思NNIE开发（三）：FasterRCNN在海思NNIE平台上的执行流程（二） 海思NNIE开发（四）:NNIE模块读入JPEG图片或视频 海思NNIE开发（五）：基...
海康AI开放平台初体验
weixin_42466834的博客
09-13 3763
..
在HI3559A上实现YOLOV3，代码可以直接用
07-27
在HI3559A上实现YOLOV3，代码可以直接用，包括输入输出和NNIE对视频文件解析。代码已测试，可以直接使用
海思3559实现yolov3实时识别
06-18
利用海思HI3559平台，来实现yolov3实时的输入输出识别处理，经过实际测试，海思3559平台可以达到8帧/s的识别速度。 本课程详细介绍了该算法实现的原理，框架以及过程，详细解释了源代码，并且已经将实现的代码上传课件，学员可以下载后直接实现该算法。
海思3559使用Yolov3进行目标检测
dx0014的博客
03-12 3086
1、实验环境 Ubuntu虚拟机+一根网线+USB转串口+小相机（支持IMX334等，可以查看） 2、资料准备 《HiMPP V4.0 媒体处理软件开发参考.pdf》 这个资料主要介绍了图像是通过什么方式进行输入，处理，以及输出的。 VI-VPSS-VO，VENC和VDEC，VGS画图等 《HiSVP开发指南.pdf》 这个资料主要介绍如何使用NNIE模块进行图像处理，以及ruyistudio的安装和使用教程。 《HiSVP API 参考.pdf》等 API等资料可以编写代码的时候在重点关注。 推荐想学习的
【项目--Hi3559A】如何在Hi3559A上运行自己的yolov3模型（修改类别、网络结构）
青玉白露的博客
07-22 2032
一、前言 使用的SDK版本是030，板子与虚拟机Linux系统之间已进行了nfs挂载。 PS：吐槽一句，这海思的sample代码写的极其不人性化，许多需要修改的参数没有写在一个地方，其中一个参数一直报段错误，使团队卡了两天······ 遇到问题不要慌，使劲撸源码就对了，“华为海思天下第一” 二、Sample代码调用分析 三、修改相关参数 四、调试脚本分享 （还在施工，未完待续） ...
海思3559A yolov3仿真代码自搭建
luozhichengaichenlei的博客
04-14 862
0、目的 是为了后面比较方便修改代码到yolov3-tiny，自搭建一份代码就不会破坏海思原始的工程。 1、搭建 首先是找到所有.h和.cpp，这里因为代码里面涉及到ssd、yolov1、yolov2，不好分离出来，还是一并放进来，只是main函数不调用就行，所有的文件如下图所示，所有代码均直接复制不用修改： 配置环境，这里我是opencv版的仿真： 这里我除了openc...
海思3559AV100平台的opencv4.0.1
05-13
综上所述，海思3559AV100平台的OpenCV 4.0.1是一个专为这款处理器优化的计算机视觉库，通过交叉编译确保了在目标平台上的运行效率。开发者可以利用这个版本进行图像处理、物体检测、人脸识别等计算机视觉任务，同时...
海思Hi3559AV100芯片硬件开发板评估板CADENCE（原理图+ PCB）工程文件+PDF原理图
04-25
海思Hi3559AV100是一款高性能的视频处理芯片，主要应用于安防监控、智能设备以及多媒体系统等领域。该芯片集成了多种功能，包括高清视频编码、解码、图像信号处理、智能分析等，为硬件开发提供了一个强大的平台。 ...
海思Hi3559V100 SDK应用指南
11-26
海思Hi3559V100 SDK应用指南，包含系统初始化、Ndk 集成开发、中间件集成等
海思Hi3559AV100开发板CADENCE原理图和 PADS9.5 PCB源文件.zip
01-30
海思Hi3559AV100是一款高性能的芯片，专为视频处理和智能物联网应用设计。这款芯片集成了多种功能，包括图像信号处理器（ISP）、视频编码器、解码器以及强大的中央处理器，使其在监控摄像头、智能家居、自动驾驶等...
海思3559AV100ES.rar
07-19
3. **集成GPU**：海思3559AV100ES内置图形处理单元（GPU），如Mali系列，用于处理图像和视频处理，支持高清显示，适用于多媒体应用和图形界面操作。 4. **网络连接**：作为一款物联网芯片，3559AV100ES可能集成了Wi...
HI3559学习（四）-- YOLOV3实时识别（实测8帧每秒）
阳光非宅男的博客
01-13 1万+
一 概述 在运行了demo 之后提示 Usage : ./sample_nnie_main <index> index: 0) RFCN(VI->VPSS->NNIE->VGS->VO). 1) Segnet(Read File). 2) FasterRcnnAlexnet(Read File). ...
HISI3559A 使用yolov3 （VI-VPSS-VO）实时目标检测
12-17 3416
使用海思3559A yolov3实时目标检测 VI-VPSS-VO
如何在Hi3559A上运行自己的YOLOv3模型(一)
qq_34533248的博客
10-11 3034
完成3559A的Hi交叉编译器的安装并解压sdk包。 step1. 复制并拷贝 aarch64-himix100-linux.tgz 与 Hi3559AV100_SDK_V2.0.2.0.tar 到 linux下 step2. 在拷贝后的目录下解压并安装交叉编译器，进入拷贝后的目录执行 tar -zxf aarch64-himix100-linux.tgz cd aarch64...
海思hi3559AV100编译GPU驱动
阿Sir问道
05-31 3851
1、前言 GPU （Graphics Process Unit ）是3D图形处理单元，海思提供完全符合业界标准的 EGL 1.4，OpenGL ES 1.1/2.0/3.0/3.1/3.2 及OpenCL2.0 标准接口。 2、驱动编译和加载 GPU 代码路径：SDK 目录下mpp/component/gpu文件夹。GPU 驱动分内核态驱动与用户态驱动...
如何在Hi3559A上运行自己的YOLOv3模型
qq_34533248的博客
10-11 2928
发现现在相关资源大多比较混乱，看起来比较难有头绪，入门难度比较大，所以单独整理并记录自己在3559A上的学习yolo模型的过程 此贴为索引贴，下为目录。博主正在加紧更新中.... 注：海思官方推荐的编译环境为windows + linux +Hi3559A共同完成编译，分工如下 ...
无人机集群分层通信及连通性研究
最新发布
10-18
针对强干扰环境，在采用窄波束固定安装的通信模块执行任务的过程中，采用一种无人机集群分层通信与连通性判断的方法，通过上下层无人机两两之间的通信关系，确定无人机集群的连通性。局部信息在集群中多次中转实现整体的信息共享，为绕飞搜索、队形保持、防止碰撞等措施的实现提供信息支持，是执行高风险、高难度复杂任务的一种高可靠、低成本的 方案。
海思3559av100 BSP FAQ文档：问题解决方案指南
"海思3559av100 BSP FAQ文档是针对海思半导体的一款芯片Hi3559AV100的板级支持包（BSP）的常见问题解答集合。该文档旨在解决在开发和使用过程中遇到的问题，提供解决方案。文档特别指出，若无特别说明，Hi3559CV100...
评论 8
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值