【进迭时空 K1 MUSE Pi Pro】基于OpenCL的FastDepth部署和扩展板设计

一、扩展板设计

1. 接口

12V电源输入接口
5V电压输出接口
电机三相输入接口
硬件I2C接口
ST7789/LCD 驱动接口
USB转TTL接口

2. 原理图

3. PCB

4. 实物图

二、基于OpenCL的FastDepth部署

1. python环境搭建

首先安装虚拟环境容器

sudo apt-get install python3-virtualenv

之后运行命令创建虚拟环境，与主环境进行隔离，myenv可以换成任意名字

virtualenv myenv

运行虚拟环境激活脚本，进入虚拟环境

source myenv/bin/activate

升级pip版本

pip install --upgrade pip

设置pip安装源

pip config set global.extra-index-url https://git.spacemit.com/api/v4/projects/33/packages/pypi/simple

这个项目需要用到onnx，onnxruntime，提取onnx模型参数，所以安装这几个库。如果安装opencv-python会直接安装python3.10版本的opencv-python，而python版本是3.12，会有兼容错误，这里用PIL库替代

pip install onnx
pip install onnxruntime
pip install pillow

2. c++环境搭建

项目需要用到opencv获取图像数据，以及cnpy库读取权重参数，所以克隆这两个库

git clone https://github.com/opencv/opencv.git -b 4.x --depth 1
git clone https://github.com/opencv/opencv_contrib.git -b 4.x --depth 1
git clone https://github.com/rogersce/cnpy.git -b master --depth 1

对cnpy进行编译

cd cnpy
mkdir build
cd build
cmake .. -DCMAKE_ISNTALL_PREFIX=../install
make -j4
sudo make install

对opencv进行编译

cd opencv
mkdir build
cd build
cmake .. -DCMAKE_INSTALL_PREFIX=../install -DCMAKE_BUILD_TYPE=Release -DOPENCV_EXTRA_MODULES_PATH = ../../opencv_contrib
make -j4
sudo make -install

3. ONNX参数提取

该项目选用fastdepth的onnx模型进行部署

https://github.com/JustinFincher/FastDepth-CoreML/blob/master/Models/fastestdepth.onnx

创建一个python脚本处理onnx，主要做下面的三个步骤

提取各个卷积层的权重参数
各个卷积层参数的数量
各个卷积层的padding，stride和kernel_shape
各个层的输入和输出张量的维度大小

代码如下：

import onnx
from onnx import helper, numpy_helper
import onnxruntime as ort
import numpy as np

model_path = "/home/ray/Documents/i2c_lcd/fastestdepth.onnx"

WEIGHTS_CPP_PATH = "/home/ray/Documents/i2c_lcd/fastdepth/source/weights.cpp"
WEIGHTS_NPY_PATH = "/home/ray/Documents/i2c_lcd/fastdepth/npy/weights.npy"

if __name__ == "__main__":
  model = onnx.load(model_path)

#.................... 新建weights.cpp写入参数
  f = open(WEIGHTS_CPP_PATH, "w")
  f.write("#include \"network.hpp\"\n")
#....................

# 获取各个层的输出
  for node in model.graph.node:
    for o in node.output:
      model.graph.output.extend([onnx.ValueInfoProto(name=o)])

# 进行推理
  session = ort.InferenceSession(model.SerializePartialToString())
  inputs = session.get_inputs()
  outputs = session.get_outputs()
  outputs_shape = {}
  for i in range(len(outputs)):
    outputs_shape[outputs[i].name] = outputs[i].shape

  onnx_weights = {}
  for initializer in model.graph.initializer:
    onnx_weights[initializer.name] = numpy_helper.to_array(initializer)
  
  new_onnx_weights = []
  weight_names = onnx_weights.keys()
  parameters_num = 0
  nodes_num = 0

  weights = {}
  weights_pad_stride = {}
  cnt = 0
  for node in model.graph.node:
    name = node.name
    attr = node.attribute
    out = node.output[0]

    for ins in node.input:
      if ins in weight_names:
        new_name = ins
        nodes_num += 1
        flatten_weight = onnx_weights[ins].flatten()
        n = len(flatten_weight)
        if len(onnx_weights[ins].shape) == 4:
          d = {"pads": None, "strides": None}
          for a in attr:
            if a.name == "group":
              if a.i != onnx_weights[ins].shape[0]:
                new_name += "Conv"
              else:
                new_name += "DepthWise"
            elif a.name == "pads":
              d["pads"] = a.ints
            elif a.name == "strides":
              d["strides"] = a.ints
            elif a.name == "kernel_shape":
              d["kernel_shape"] = a.ints
            d["output_c"] = outputs_shape[out][1]
          weights_pad_stride[new_name] = d
        elif len(onnx_weights[ins].shape) == 1:
          new_name += "Bias"
        
        weights[new_name] = [parameters_num, n]

        parameters_num += n

        for w in flatten_weight:
          new_onnx_weights.append(w)

  f.write("\nvoid Network::init_weights(const std::string & npy_path)\n")
  f.write("{\n")
  f.write("  parameters_num = {};\n".format(parameters_num))
  f.write("  npz_array = cnpy::npy_load(npy_path);\n")
  f.write("  weights = npz_array.data<float>();\n\n")
  for w in weights:
    f.write("  pos[\"{}\"].start = {}; pos[\"{}\"].num = {};\n".format(w, weights[w][0], w, weights[w][1]))
  f.write("\n")

  for w in weights_pad_stride:
    f.write("  weights_param[\"{}\"].padding = {}; weights_param[\"{}\"].stride = {}; weights_param[\"{}\"].kernel_shape = {}; weights_param[\"{}\"].out_channel = {};".format(
      w, weights_pad_stride[w]["pads"][0], w, weights_pad_stride[w]["strides"][0],
      w, weights_pad_stride[w]["kernel_shape"][0], w, weights_pad_stride[w]["output_c"]
    ))
    f.write("\n")
  f.write("\n")

  np.save(WEIGHTS_NPY_PATH, new_onnx_weights)

  image = np.random.uniform(0, 1, size=(1, inputs[0].shape[1], inputs[0].shape[2], inputs[0].shape[3])).astype(np.float32)
  inference_outputs = session.run(None, {inputs[0].name: image})

  inference_outputs_names = [x.name for x in session.get_outputs()]
  inference_outputs_dict = dict()
  for i in range(len(inference_outputs_names)):
    key = inference_outputs_names[i]
    value = inference_outputs[i]
    inference_outputs_dict[key] = value

#................
    f.write("  inout_shape[\"{}\"][0] = {}; ".format(key, value.shape[0]))
    f.write("inout_shape[\"{}\"][1] = {};".format(key, value.shape[1]))
    f.write("inout_shape[\"{}\"][2] = {};".format(key, value.shape[2]))
    f.write("inout_shape[\"{}\"][3] = {};\n".format(key, value.shape[3]))
#................

  for x in session.get_inputs():
    f.write("  inout_shape[\"{}\"][0] = {}; ".format(x.name, x.shape[0]))
    f.write("inout_shape[\"{}\"][1] = {};".format(x.name, x.shape[1]))
    f.write("inout_shape[\"{}\"][2] = {};".format(x.name, x.shape[2]))
    f.write("inout_shape[\"{}\"][3] = {};\n".format(x.name, x.shape[3]))

  f.write("} \n")

执行上述代码后生成weights.cpp和weights.npy文件，其中weights.npy包含所有权重参数，weights.cpp记录各个参数获取的起始索引和数量、输入输出维度和卷积层的padding等参数

3. 基于OpenCL的算子开发

fastdepth主要包含5种算子操作

二维卷积
可分离卷积
ReLU激活
张量相加
上采样

核函数设计如下

二维卷积

const char * kernel_conv2d = KERNEL(
__kernel void conv2d(
  __global const float * input,
  __global const float * weights,
  __global const float * bias,
  __global float * output,
  __private int kernel_chnlnum,
  __private int kernel_hw,
  __private int stride,
  __private int padding,
  __private int input_height,
  __private int input_width,
  __private float clip_min,
  __private float clip_max
)
{
  const int c = get_global_id(0);
  const int h = get_global_id(1);
  const int w = get_global_id(2);

  const int output_height = get_global_size(1);
  const int output_width = get_global_size(2);

  float sum = 0;
  int woffset = c * kernel_chnlnum * kernel_hw * kernel_hw;

  int hidx = kernel_hw / 2 - padding + h * stride;
  int widx = kernel_hw / 2 - padding + w * stride;

  for (int i = 0; i < kernel_chnlnum; i++)
  {
    for (int j = 0; j < kernel_hw; j++)
    {
      int input_hidx = hidx + j - kernel_hw / 2;
      if ((input_hidx < 0) || (input_hidx >= input_height))
      {
        continue;
      }
      for (int k = 0; k < kernel_hw; k++)
      {
        int input_widx = widx + k - kernel_hw / 2;
        if ((input_widx < 0) || (input_widx >= input_width))
        {
          continue;
        }

        float weight = weights[woffset + i * kernel_hw * kernel_hw + j * kernel_hw + k];
        float inp = input[i * input_height * input_width + input_hidx * input_width + input_widx];

        sum += weight * inp;
      }
    }
  }

  sum += bias[c];

  if (clip_min < clip_max)
  {
    if (sum < clip_min)
    {
      sum = clip_min;
    }
    else if (sum > clip_max)
    {
      sum = clip_max;
    }
  }

  output[c * output_width * output_height + h * output_width + w] = sum;
}
);

可分离卷积

const char * kernel_depthwise_conv2d = KERNEL(
__kernel void depthwise_conv2d(
  __global const float * input,
  __global const float * weights,
  __global const float * bias,
  __global float * output,
  __private int kernel_hw,
  __private int stride,
  __private int padding,
  __private int input_height,
  __private int input_width,
  __private float clip_min,
  __private float clip_max
)
{
  const int c = get_global_id(0);
  const int h = get_global_id(1);
  const int w = get_global_id(2);

  const int output_height = get_global_size(1);
  const int output_width = get_global_size(2);

  float sum = 0;
  int woffset = c * kernel_hw * kernel_hw;
  int ioffset = c * input_height * input_width;

  int hidx = kernel_hw / 2 - padding + h * stride;
  int widx = kernel_hw / 2 - padding + w * stride;

  for (int j = 0; j < kernel_hw; j++)
  {
    int input_hidx = hidx + j - kernel_hw / 2;
    if ((input_hidx < 0) || (input_hidx >= input_height))
    {
      continue;
    }
    for (int k = 0; k < kernel_hw; k++)
    {
      int input_widx = widx + k - kernel_hw / 2;
      if ((input_widx < 0) || (input_widx >= input_width))
      {
        continue;
      }  
      float weight = weights[woffset + j * kernel_hw + k];
      float inp = input[ioffset + input_hidx * input_width + input_widx];  
      sum += weight * inp;
    }
  }

  sum += bias[c];

  if (clip_min < clip_max)
  {
    if (sum < clip_min)
    {
      sum = clip_min;
    }
    else if (sum > clip_max)
    {
      sum = clip_max;
    }
  }

  output[c * output_width * output_height + h * output_width + w] = sum;
}
);

张量相加

const char * kernel_add = KERNEL(
__kernel void add(
  __global const float * input1,
  __global const float * input2,
  __global float * output
)
{
  const int c = get_global_id(0);
  const int h = get_global_id(1);
  const int w = get_global_id(2);

  const int h_size = get_global_size(1);
  const int w_size = get_global_size(2);

  const int tid = c * h_size * w_size + h * w_size + w;

  output[tid] = input1[tid] + input2[tid];
}
);

ReLU激活

const char * kernel_relu = KERNEL(
__kernel void relu(
  __global const float * input,
  __global float * output
)
{
  const int c = get_global_id(0);
  const int h = get_global_id(1);
  const int w = get_global_id(2);

  const int h_size = get_global_size(1);
  const int w_size = get_global_size(2);

  const int tid = c * h_size * w_size + h * w_size + w;

  if (input[tid] > 0)
  {
    output[tid] = input[tid];
  }
  else
  {
    output[tid] = 0.0;
  }
}
);

上采样（nearest方法）

const char * kernel_upsample_nearest = KERNEL(
__kernel void upsample_nearest(
  __global const float * input,
  __global float * output,
  __private int input_height,
  __private int input_width,
  __private int h_factor,
  __private int w_factor
)
{
  const int c = get_global_id(0);
  const int h = get_global_id(1);
  const int w = get_global_id(2);

  const int h_size = get_global_size(1);
  const int w_size = get_global_size(2);

  const int in_h = h / h_factor;
  const int in_w = w / w_factor;

  const int out_idx = c * h_size * w_size + h * w_size + w;
  const int in_idx = c * input_height * input_width + in_h * input_width + in_w;

  output[out_idx] = input[in_idx];
}
);

4. FastDepth算法部署

在network.cpp里提供创建卷积层、可分离卷积层等的接口

卷积层

void Network::depthwise_conv2d_layer(const std::string & input_name,
                            const std::string & output_name,
                            const std::string & weight_name,
                            const std::string & bias_name,
                            const char * kernel_name,
                            const float & clip_min,
                            const float & clip_max)
{
  cl_int errcode;

  if (cl_nodes.find(input_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input_name][1] * inout_shape[input_name][2] * inout_shape[input_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(weight_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[weight_name] = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(float) * pos[weight_name].num, &weights[ pos[weight_name].start ], &errcode));
  }
  if (cl_nodes.find(bias_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[bias_name] = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(float) * pos[bias_name].num, &weights[ pos[bias_name].start ], &errcode));
  }
  if (cl_nodes.find(output_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[output_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[output_name][1] * inout_shape[output_name][2] * inout_shape[output_name][3], NULL, &errcode));
  }

  if (cl_kernels.find(kernel_name) == cl_kernels.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_kernels[kernel_name] = clCreateKernel(cl_programs["depthwise_conv2d"], "depthwise_conv2d", &errcode));
  }

  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 0, sizeof(cl_mem), &cl_nodes[input_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 1, sizeof(cl_mem), &cl_nodes[weight_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 2, sizeof(cl_mem), &cl_nodes[bias_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 3, sizeof(cl_mem), &cl_nodes[output_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 4, sizeof(int), &weights_param[weight_name].kernel_shape));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 5, sizeof(int), &weights_param[weight_name].stride));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 6, sizeof(int), &weights_param[weight_name].padding));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 7, sizeof(int), &inout_shape[input_name][2]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 8, sizeof(int), &inout_shape[input_name][3]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 9, sizeof(float), &clip_min));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 10, sizeof(float), &clip_max));

  kernel_size[kernel_name][0] = inout_shape[output_name][1];
  kernel_size[kernel_name][1] = inout_shape[output_name][2];
  kernel_size[kernel_name][2] = inout_shape[output_name][3];
}

可分离卷积层

void Network::depthwise_conv2d_layer(const std::string & input_name,
                            const std::string & output_name,
                            const std::string & weight_name,
                            const std::string & bias_name,
                            const char * kernel_name,
                            const float & clip_min,
                            const float & clip_max)
{
  cl_int errcode;

  if (cl_nodes.find(input_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input_name][1] * inout_shape[input_name][2] * inout_shape[input_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(weight_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[weight_name] = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(float) * pos[weight_name].num, &weights[ pos[weight_name].start ], &errcode));
  }
  if (cl_nodes.find(bias_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[bias_name] = clCreateBuffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, sizeof(float) * pos[bias_name].num, &weights[ pos[bias_name].start ], &errcode));
  }
  if (cl_nodes.find(output_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[output_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[output_name][1] * inout_shape[output_name][2] * inout_shape[output_name][3], NULL, &errcode));
  }

  if (cl_kernels.find(kernel_name) == cl_kernels.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_kernels[kernel_name] = clCreateKernel(cl_programs["depthwise_conv2d"], "depthwise_conv2d", &errcode));
  }

  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 0, sizeof(cl_mem), &cl_nodes[input_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 1, sizeof(cl_mem), &cl_nodes[weight_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 2, sizeof(cl_mem), &cl_nodes[bias_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 3, sizeof(cl_mem), &cl_nodes[output_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 4, sizeof(int), &weights_param[weight_name].kernel_shape));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 5, sizeof(int), &weights_param[weight_name].stride));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 6, sizeof(int), &weights_param[weight_name].padding));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 7, sizeof(int), &inout_shape[input_name][2]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 8, sizeof(int), &inout_shape[input_name][3]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 9, sizeof(float), &clip_min));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 10, sizeof(float), &clip_max));

  kernel_size[kernel_name][0] = inout_shape[output_name][1];
  kernel_size[kernel_name][1] = inout_shape[output_name][2];
  kernel_size[kernel_name][2] = inout_shape[output_name][3];
}

relu激活

void Network::relu_layer(const std::string & input_name,
                const std::string & output_name,
                const char * kernel_name)
{
  cl_int errcode;

  if (cl_nodes.find(input_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input_name][1] * inout_shape[input_name][2] * inout_shape[input_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(output_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[output_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[output_name][1] * inout_shape[output_name][2] * inout_shape[output_name][3], NULL, &errcode));
  }

  if (cl_kernels.find(kernel_name) == cl_kernels.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_kernels[kernel_name] = clCreateKernel(cl_programs["relu"], "relu", &errcode));
  }

  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 0, sizeof(cl_mem), &cl_nodes[input_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 1, sizeof(cl_mem), &cl_nodes[output_name]));

  kernel_size[kernel_name][0] = inout_shape[output_name][1];
  kernel_size[kernel_name][1] = inout_shape[output_name][2];
  kernel_size[kernel_name][2] = inout_shape[output_name][3];
}

张量相加

void Network::add_layer(const std::string & input1_name,
               const std::string & input2_name,
               const std::string & output_name,
               const char * kernel_name)
{
  cl_int errcode;

  if (cl_nodes.find(input1_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input1_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input1_name][1] * inout_shape[input1_name][2] * inout_shape[input1_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(input2_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input2_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input2_name][1] * inout_shape[input2_name][2] * inout_shape[input2_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(output_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[output_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[output_name][1] * inout_shape[output_name][2] * inout_shape[output_name][3], NULL, &errcode));
  }

  if (cl_kernels.find(kernel_name) == cl_kernels.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_kernels[kernel_name] = clCreateKernel(cl_programs["add"], "add", &errcode));
  }

  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 0, sizeof(cl_mem), &cl_nodes[input1_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 1, sizeof(cl_mem), &cl_nodes[input2_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 2, sizeof(cl_mem), &cl_nodes[output_name]));

  kernel_size[kernel_name][0] = inout_shape[output_name][1];
  kernel_size[kernel_name][1] = inout_shape[output_name][2];
  kernel_size[kernel_name][2] = inout_shape[output_name][3];
}

上采样层

void Network::upsample_nearest_layer(const std::string & input_name,
                            const std::string & output_name,
                            const char * kernel_name,
                            const int & h_factor,
                            const int & w_factor)
{
  cl_int errcode;

  if (cl_nodes.find(input_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[input_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[input_name][1] * inout_shape[input_name][2] * inout_shape[input_name][3], NULL, &errcode));
  }
  if (cl_nodes.find(output_name) == cl_nodes.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_nodes[output_name] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * inout_shape[output_name][1] * inout_shape[output_name][2] * inout_shape[output_name][3], NULL, &errcode));
  }

  if (cl_kernels.find(kernel_name) == cl_kernels.end())
  {
    CHECK_ERROR(errcode != CL_SUCCESS, errcode, cl_kernels[kernel_name] = clCreateKernel(cl_programs["upsample_nearest"], "upsample_nearest", &errcode));
  }

  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 0, sizeof(cl_mem), &cl_nodes[input_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 1, sizeof(cl_mem), &cl_nodes[output_name]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 2, sizeof(int), &inout_shape[input_name][2]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 3, sizeof(int), &inout_shape[input_name][3]));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 4, sizeof(int), &h_factor));
  CHECK_ERROR(errcode != CL_SUCCESS, errcode, errcode = clSetKernelArg(cl_kernels[kernel_name], 5, sizeof(int), &w_factor));

  kernel_size[kernel_name][0] = inout_shape[output_name][1];
  kernel_size[kernel_name][1] = inout_shape[output_name][2];
  kernel_size[kernel_name][2] = inout_shape[output_name][3];
}

main.cpp构建fastdpeth结构

int main(int argc, char ** argv)
{
  std::cout << "starting..." << std::endl;

  try
  {
    std::string save_npy_path;
    if (argc > 1)
    {
      save_npy_path = std::string(argv[1]);
    }

    Network net;
    net.init_weights("/home/ray/Documents/i2c_lcd/fastdepth/npy/weights.npy");
    net.init();

    net.init_input_layer("input", "data", "init_input");
    net.conv2d_layer("data", "231", "507Conv", "509Bias", "conv2d_507", 0, 6);
    net.depthwise_conv2d_layer("231", "234", "511DepthWise", "513Bias", "depthwise_conv2d_511", 0, 6);
    net.conv2d_layer("234", "237", "515Conv", "517Bias", "conv2d_515", 0, 6);
    net.depthwise_conv2d_layer("237", "240", "519DepthWise", "521Bias", "depthwise_conv2d_519", 0, 6);
    net.conv2d_layer("240", "243", "523Conv", "525Bias", "conv2d_523", 0, 6);
    net.depthwise_conv2d_layer("243", "246", "527DepthWise", "529Bias", "depthwise_conv2d_527", 0, 6);
    /*
      ..............
    */

    net.execute(save_npy_path);
    net.release();
  }
  catch(const std::exception & e)
  {
    std::cerr << e.what() << std::endl;
  }

  std::cout << "输入任意字符结束..." << std::endl;
  std::cin.get();

  return 0;
}

5. 效果展示

视差图生成效果

运行速度效果：

【进迭时空 K1 MUSE Pi Pro】基于OpenCL的FastDepth部署和扩展板设计

PCBA

设计文件

生产BOM

源代码

详细介绍

【进迭时空 K1 MUSE Pi Pro】基于OpenCL的FastDepth部署和扩展板设计

一、扩展板设计

1. 接口

2. 原理图

3. PCB

4. 实物图

二、基于OpenCL的FastDepth部署

1. python环境搭建

2. c++环境搭建

3. ONNX参数提取

3. 基于OpenCL的算子开发

4. FastDepth算法部署

5. 效果展示