CUDA与TensorRT | 把检测器加进来,YOLOv8部署实战!

2023-12-16

点击下方 卡片 ,关注“ 自动驾驶之心 ”公众号

ADAS巨卷干货,即可获取

>> 点击进入→ 自动驾驶之心【模型部署】技术交流群

论文作者 | 汽车人

编辑 | 自动驾驶之心

0. 把检测器加进来

本文是我在学习韩博《CUDA与TensorRT部署实战课程》第六章的课程部分输出的个人学习笔记,欢迎大家一起讨论学习!

1. 导出onnx需要注意的地方

不要pip install ultralytics 而是选择git clone 的方式安装yolov8

# Clone the ultralytics repository
git clone https://github.com/ultralytics/ultralytics

# Navigate to the cloned directory
cd ultralytics

# Install the package in editable mode for development
pip install -e .

然后导出onnx看一下

model = YOLO('yolov8n.pt')  # load an official model
model = YOLO('path/to/best.pt')  # load a custom trained model

# Export the model
model.export(format='onnx')
4ab25578b711f4e9189b4e35cbd6e0ae.png

这个onnx长这样,但是为了速度能够更好一些,因为内存是连续的,为了提升速度我们需要把这个输出变成1x8400x84, 因为这样我们的数据才是8400 个bbox信息(cx, cy, w, h)连在一起, 不然根据官方的信息这里是每8400个cx, 然后是每8400个cy这样的,所以这里需要对他进行一个转置,简单的代码如下:

# ultralytics/ultralytics/nn/modules/head.py
class Detect(nn.Module):
    # ...
    def forward(self, x):
        # ...
        y = torch.cat((dbox, cls.sigmoid()), 1)
        y = y.transpose(1, 2)
        return y if self.export else (y, x)

这里加了一个转置之后,他就会变成

3b50396a39389b48cf45f23317710ab8.png

这样我们的yolov8的onnx就成功的导出了

2d35a1f473e42e69a7f0acd520d1d4b1.png

a66949511bb393f18945cc776d788549.png

扫码领取八折优惠券!

2. 看一下在分类器的框架上这里修改了哪些地方

这里再复习一下这个推理框架的操作

首先, Worker 是一个类,它负责管理模型的生命周期,包括模型的初始化、加载图像、进行推理等。它根据任务类型(分类或检测)来创建相应的模型实例。

然后,我们有两种类型的模型:Classifier 和 Detector 。这两种模型都继承自 Model 基类。

Classifier 是用于图像分类任务的模型。它包含了一些特定于分类任务的方法,如预处理和后处理。

Detector 是用于目标检测任务的模型。它也包含了一些特定于检测任务的方法,如预处理和后处理。

这两种模型都有自己的 setup 方法来初始化模型,包括创建推理引擎、设置输入/输出绑定、分配内存等。此外,它们还有自己的 preprocess 和 postprocess 方法来处理输入图像和输出结果。

3b01d8d0bf41a0096aeb908a2b97f098.png

主要更改下面几个文件

  • src/cpp/trt_detector.cpp

  • src/cpp/trt_worker.cpp

  • src/cpp/trt_model.cpp

  • include/trt_detector.hpp

  • inlcude/trt_worker.hpp

下面是整个框架的一个流程图

f39181324c5508c5ef8e4da1c26963fd.png

其实通过上面的流程图可以简单的看出来除了前后处理其他的都是差不多的, 所以也就是前面worker加了一个m_detector, 然后加了一些bbox的数据结构,

所以这里重点看这个前后处理

3. 前处理

读取图片, 然后调用一下前处理函数, 这里的以GPU版本的前处理为例子, 下面是在Detector类实现的, 先读取图片, 然后从param读取信息, 然后这里调用

bool Detector::preprocess_gpu() {
    /*Preprocess -- yolo的预处理并没有mean和std,所以可以直接skip掉mean和std的计算 */

    /*Preprocess -- 读取数据*/
    m_inputImage = cv::imread(m_imagePath);
    if (m_inputImage.data == nullptr) {
        LOGE("ERROR: file not founded! Program terminated"); return false;
    }
    
    /*Preprocess -- 测速*/
    m_timer->start_gpu();

    /*Preprocess -- 使用GPU进行warpAffine, 并将结果返回到m_inputMemory中*/
    preprocess::preprocess_resize_gpu(m_inputImage, m_inputMemory[1],
                                   m_params->img.h, m_params->img.w, 
                                   preprocess::tactics::GPU_WARP_AFFINE);

    m_timer->stop_gpu();
    m_timer->duration_gpu("preprocess(GPU)");
    return true;
}

这里就是正常的核函数的一个步骤, 通过分配内存然后把数据从Host搬到Device然后开始执行cu里面的函数, 再借由执行cu文件里面的函数调用kernel函数

void preprocess_resize_gpu(
    cv::Mat &h_src, float* d_tar, 
    const int& tar_h, const int& tar_w, 
    tactics tac) 
{
    uint8_t* d_src  = nullptr;

    int height   = h_src.rows;
    int width    = h_src.cols;
    int chan     = 3;

    int src_size  = height * width * chan * sizeof(uint8_t);
    int norm_size = 3 * sizeof(float);


    // 分配device上的src的内存
    CUDA_CHECK(cudaMalloc(&d_src, src_size));

    // 将数据拷贝到device上
    CUDA_CHECK(cudaMemcpy(d_src, h_src.data, src_size, cudaMemcpyHostToDevice));

    // device上处理resize, BGR2RGB的核函数
    resize_bilinear_gpu(d_tar, d_src, tar_w, tar_h, width, height, tac);

    // host和device进行同步处理
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaFree(d_src));

    // 因为接下来会继续在gpu上进行处理,所以这里不用把结果返回到host
}

这个是cu里面的核函数, 这里主要是用来调用核函数,这里有很多种, 通过选择的方法来进行不同的核函数:

warpaffine_init(srcH, srcW, tarH, tarW);
        warpaffine_BGR2RGB_kernel <<<dimGrid, dimBlock>>> 
                (d_tar, d_src, trans, affine_matrix);
void resize_bilinear_gpu(
    float* d_tar, uint8_t* d_src, 
    int tarW, int tarH, 
    int srcW, int srcH, 
    tactics tac) 
{
    dim3 dimBlock(32, 32, 1);
    dim3 dimGrid(tarW / 32 + 1, tarH / 32 + 1, 1);
   
    //scaled resize
    float scaled_h = (float)srcH / tarH;
    float scaled_w = (float)srcW / tarW;
    float scale = (scaled_h > scaled_w ? scaled_h : scaled_w);

    switch (tac) {
    case tactics::GPU_NEAREST:
        nearest_BGR2RGB_nhwc2nchw_kernel <<<dimGrid, dimBlock>>>
                (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case tactics::GPU_NEAREST_CENTER:
        nearest_BGR2RGB_nhwc2nchw_kernel <<<dimGrid, dimBlock>>>
                (d_tar, d_src, tarW, tarH, srcW, srcH, scale, scale);
        break;
    case tactics::GPU_BILINEAR:
        bilinear_BGR2RGB_nhwc2nchw_kernel <<<dimGrid, dimBlock>>> 
                (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case tactics::GPU_BILINEAR_CENTER:
        bilinear_BGR2RGB_nhwc2nchw_shift_kernel <<<dimGrid, dimBlock>>> 
                (d_tar, d_src, tarW, tarH, srcW, srcH, scale, scale);
        break;
    case tactics::GPU_WARP_AFFINE:
        warpaffine_init(srcH, srcW, tarH, tarW);
        warpaffine_BGR2RGB_kernel <<<dimGrid, dimBlock>>> 
                (d_tar, d_src, trans, affine_matrix);
        break;
    default:
        LOGE("ERROR: Wrong GPU resize tactics selected. Program terminated");
        exit(1);
    }
}

这里是对这些函数的解释:

  1. warpaffine_init :

    void warpaffine_init(int srcH, int srcW, int tarH, int tarW){
     trans.src_h = srcH;
     trans.src_w = srcW;
     trans.tar_h = tarH;
     trans.tar_w = tarW;
     affine_matrix.init(trans);
    }
  • 初始化仿射变换所需的参数。它设置了源图像和目标图像的高度和宽度,并初始化一个 AffineMatrix 对象。

affine_transformation :

  • 这是一个在CPU和GPU上都可以调用的函数,用于执行仿射变换。它根据传入的变换矩阵和源坐标计算目标坐标。

__host__ __device__ void affine_transformation(
    float trans_matrix[6], 
    int src_x, int src_y, 
    float* tar_x, float* tar_y)
{
    *tar_x = trans_matrix[0] * src_x + trans_matrix[1] * src_y + trans_matrix[2];
    *tar_y = trans_matrix[3] * src_x + trans_matrix[4] * src_y + trans_matrix[5];
}
  1. nearest_BGR2RGB_nhwc2nchw_norm_kernel :

  • 这是一个CUDA核函数,用于将图像从BGR格式转换为RGB格式,同时将图像从NHWC格式转换为NCHW格式,并应用最邻近插值方法进行缩放。它还包括对像素值进行归一化处理(减去均值,除以标准差)。

__global__ void nearest_BGR2RGB_nhwc2nchw_norm_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH,
    float scaled_w, float scaled_h,
    float* d_mean, float* d_std) 
{
    // nearest neighbour -- resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // nearest neighbour -- 计算最近坐标
    int src_y = floor((float)y * scaled_h);
    int src_x = floor((float)x * scaled_w);

    if (src_x < 0 || src_y < 0 || src_x > srcW || src_y > srcH) {
        // nearest neighbour -- 对于越界的部分,不进行计算
    } else {
        // nearest neighbour -- 计算tar中对应坐标的索引
        int tarIdx  = y * tarW + x;
        int tarArea = tarW * tarH;

        // nearest neighbour -- 计算src中最近邻坐标的索引
        int srcIdx = (src_y * srcW + src_x) * 3;

        // nearest neighbour -- 实现nearest beighbour的resize + BGR2RGB + nhwc2nchw + norm
        tar[tarIdx + tarArea * 0] = (src[srcIdx + 2] / 255.0f - d_mean[2]) / d_std[2];
        tar[tarIdx + tarArea * 1] = (src[srcIdx + 1] / 255.0f - d_mean[1]) / d_std[1];
        tar[tarIdx + tarArea * 2] = (src[srcIdx + 0] / 255.0f - d_mean[0]) / d_std[0];
    }
}
  1. bilinear_BGR2RGB_nhwc2nchw_norm_kernel :

  • 类似于上一个函数,但使用双线性插值方法进行缩放。它也执行BGR到RGB的颜色转换、NHWC到NCHW的格式转换,并对像素值进行归一化。

__global__ void bilinear_BGR2RGB_nhwc2nchw_norm_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH, 
    float scaled_w, float scaled_h,
    float* d_mean, float* d_std) 
{

    // bilinear interpolation -- resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // // bilinear interpolation -- 计算x,y映射到原图时最近的4个坐标
    int src_y1 = floor((y + 0.5) * scaled_h - 0.5);
    int src_x1 = floor((x + 0.5) * scaled_w - 0.5);
    int src_y2 = src_y1 + 1;
    int src_x2 = src_x1 + 1;

    if (src_y1 < 0 || src_x1 < 0 || src_y2 > srcH || src_x2 > srcW) {
        // bilinear interpolation -- 对于越界的坐标不进行计算
    } else {
        // bilinear interpolation -- 计算原图上的坐标(浮点类型)在0~1之间的值
        float th   = ((y + 0.5) * scaled_h - 0.5) - src_y1;
        float tw   = ((x + 0.5) * scaled_w - 0.5) - src_x1;

        // bilinear interpolation -- 计算面积(这里建议自己手画一张图来理解一下)
        float a1_1 = (1.0 - tw) * (1.0 - th);  //右下
        float a1_2 = tw * (1.0 - th);          //左下
        float a2_1 = (1.0 - tw) * th;          //右上
        float a2_2 = tw * th;                  //左上

        // bilinear interpolation -- 计算4个坐标所对应的索引
        int srcIdx1_1 = (src_y1 * srcW + src_x1) * 3;  //左上
        int srcIdx1_2 = (src_y1 * srcW + src_x2) * 3;  //右上
        int srcIdx2_1 = (src_y2 * srcW + src_x1) * 3;  //左下
        int srcIdx2_2 = (src_y2 * srcW + src_x2) * 3;  //右下

        // bilinear interpolation -- 计算resized之后的图的索引
        int tarIdx    = y * tarW  + x;
        int tarArea   = tarW * tarH;

        // bilinear interpolation -- 实现bilinear interpolation的resize + BGR2RGB + NHWC2NCHW normalization
        // 注意,这里tar和src进行遍历的方式是不一样的
        tar[tarIdx + tarArea * 0] = 
            (round((a1_1 * src[srcIdx1_1 + 2] + 
                   a1_2 * src[srcIdx1_2 + 2] +
                   a2_1 * src[srcIdx2_1 + 2] +
                   a2_2 * src[srcIdx2_2 + 2])) / 255.0f - d_mean[2]) / d_std[2];

        tar[tarIdx + tarArea * 1] = 
            (round((a1_1 * src[srcIdx1_1 + 1] + 
                   a1_2 * src[srcIdx1_2 + 1] +
                   a2_1 * src[srcIdx2_1 + 1] +
                   a2_2 * src[srcIdx2_2 + 1])) / 255.0f - d_mean[1]) / d_std[1];

        tar[tarIdx + tarArea * 2] = 
            (round((a1_1 * src[srcIdx1_1 + 0] + 
                   a1_2 * src[srcIdx1_2 + 0] +
                   a2_1 * src[srcIdx2_1 + 0] +
                   a2_2 * src[srcIdx2_2 + 0])) / 255.0f - d_mean[0]) / d_std[0];

    }
}
  1. bilinear_BGR2RGB_nhwc2nchw_shift_norm_kernel :

  • 这个函数也执行双线性插值缩放、BGR到RGB颜色转换、格式转换,但在缩放过程中还考虑了坐标的平移(shift),并进行了像素值归一化。

__global__ void bilinear_BGR2RGB_nhwc2nchw_shift_norm_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH, 
    float scaled_w, float scaled_h,
    float* d_mean, float* d_std) 
{
    // resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // bilinear interpolation -- 计算x,y映射到原图时最近的4个坐标
    int src_y1 = floor((y + 0.5) * scaled_h - 0.5);
    int src_x1 = floor((x + 0.5) * scaled_w - 0.5);
    int src_y2 = src_y1 + 1;
    int src_x2 = src_x1 + 1;

    if (src_y1 < 0 || src_x1 < 0 || src_y2 > srcH || src_x2 > srcW) {
        // bilinear interpolation -- 对于越界的坐标不进行计算
    } else {
        // bilinear interpolation -- 计算原图上的坐标(浮点类型)在0~1之间的值
        float th   = (float)y * scaled_h - src_y1;
        float tw   = (float)x * scaled_w - src_x1;

        // bilinear interpolation -- 计算面积(这里建议自己手画一张图来理解一下)
        float a1_1 = (1.0 - tw) * (1.0 - th);  // 右下
        float a1_2 = tw * (1.0 - th);          // 左下
        float a2_1 = (1.0 - tw) * th;          // 右上
        float a2_2 = tw * th;                  // 左上

        // bilinear interpolation -- 计算4个坐标所对应的索引
        int srcIdx1_1 = (src_y1 * srcW + src_x1) * 3;  // 左上
        int srcIdx1_2 = (src_y1 * srcW + src_x2) * 3;  // 右上
        int srcIdx2_1 = (src_y2 * srcW + src_x1) * 3;  // 左下
        int srcIdx2_2 = (src_y2 * srcW + src_x2) * 3;  // 右下

        // bilinear interpolation -- 计算原图在目标图中的x, y方向上的偏移量
        y = y - int(srcH / (scaled_h * 2)) + int(tarH / 2);
        x = x - int(srcW / (scaled_w * 2)) + int(tarW / 2);

        // bilinear interpolation -- 计算resized之后的图的索引
        int tarIdx    = (y * tarW  + x) * 3;
        int tarArea   = tarW * tarH;

        // bilinear interpolation -- 实现bilinear interpolation + BGR2RGB + shift + nhwc2nchw
        tar[tarIdx + tarArea * 0] = 
            (round((a1_1 * src[srcIdx1_1 + 2] + 
                   a1_2 * src[srcIdx1_2 + 2] +
                   a2_1 * src[srcIdx2_1 + 2] +
                   a2_2 * src[srcIdx2_2 + 2])) / 255.0f - d_mean[2]) / d_std[2];

        tar[tarIdx + tarArea * 1] = 
            (round((a1_1 * src[srcIdx1_1 + 1] + 
                   a1_2 * src[srcIdx1_2 + 1] +
                   a2_1 * src[srcIdx2_1 + 1] +
                   a2_2 * src[srcIdx2_2 + 1])) / 255.0f - d_mean[1]) / d_std[1];

        tar[tarIdx + tarArea * 2] = 
            (round((a1_1 * src[srcIdx1_1 + 0] + 
                   a1_2 * src[srcIdx1_2 + 0] +
                   a2_1 * src[srcIdx2_1 + 0] +
                   a2_2 * src[srcIdx2_2 + 0])) / 255.0f - d_mean[0]) / d_std[0];
    }
}
  1. nearest_BGR2RGB_nhwc2nchw_kernel bilinear_BGR2RGB_nhwc2nchw_kernel :

  • 这两个函数类似于之前描述的 _norm 版本的函数,但它们不进行像素值的归一化处理。

__global__ void nearest_BGR2RGB_nhwc2nchw_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH,
    float scaled_w, float scaled_h)
{
    // nearest neighbour -- resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // nearest neighbour -- 计算最近坐标
    int src_y = floor((float)y * scaled_h);
    int src_x = floor((float)x * scaled_w);

    if (src_x < 0 || src_y < 0 || src_x > srcW || src_y > srcH) {
        // nearest neighbour -- 对于越界的部分,不进行计算
    } else {
        // nearest neighbour -- 计算tar中对应坐标的索引
        int tarIdx  = y * tarW + x;
        int tarArea = tarW * tarH;

        // nearest neighbour -- 计算src中最近邻坐标的索引
        int srcIdx = (src_y * srcW + src_x) * 3;

        // nearest neighbour -- 实现nearest beighbour的resize + BGR2RGB + nhwc2nchw + norm
        tar[tarIdx + tarArea * 0] = src[srcIdx + 2] / 255.0f;
        tar[tarIdx + tarArea * 1] = src[srcIdx + 1] / 255.0f;
        tar[tarIdx + tarArea * 2] = src[srcIdx + 0] / 255.0f;
    }
}
  1. bilinear_BGR2RGB_nhwc2nchw_shift_kernel :

  • 类似于 bilinear_BGR2RGB_nhwc2nchw_kernel ,但考虑了坐标的平移(shift)。

__global__ void bilinear_BGR2RGB_nhwc2nchw_shift_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH, 
    float scaled_w, float scaled_h)
{
    // resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;


    // bilinear interpolation -- 计算x,y映射到原图时最近的4个坐标
    int src_y1 = floor((y + 0.5) * scaled_h - 0.5);
    int src_x1 = floor((x + 0.5) * scaled_w - 0.5);
    int src_y2 = src_y1 + 1;
    int src_x2 = src_x1 + 1;

    if (src_y1 < 0 || src_x1 < 0 || src_y2 > srcH || src_x2 > srcW) {
        // bilinear interpolation -- 对于越界的坐标不进行计算
    } else {
        // bilinear interpolation -- 计算原图上的坐标(浮点类型)在0~1之间的值
        float th   = (float)y * scaled_h - src_y1;
        float tw   = (float)x * scaled_w - src_x1;

        // bilinear interpolation -- 计算面积(这里建议自己手画一张图来理解一下)
        float a1_1 = (1.0 - tw) * (1.0 - th);  // 右下
        float a1_2 = tw * (1.0 - th);          // 左下
        float a2_1 = (1.0 - tw) * th;          // 右上
        float a2_2 = tw * th;                  // 左上

        // bilinear interpolation -- 计算4个坐标所对应的索引
        int srcIdx1_1 = (src_y1 * srcW + src_x1) * 3;  // 左上
        int srcIdx1_2 = (src_y1 * srcW + src_x2) * 3;  // 右上
        int srcIdx2_1 = (src_y2 * srcW + src_x1) * 3;  // 左下
        int srcIdx2_2 = (src_y2 * srcW + src_x2) * 3;  // 右下

        // bilinear interpolation -- 计算原图在目标图中的x, y方向上的偏移量
        y = y - int(srcH / (scaled_h * 2)) + int(tarH / 2);
        x = x - int(srcW / (scaled_w * 2)) + int(tarW / 2);

        // bilinear interpolation -- 计算resized之后的图的索引
        int tarIdx    = y * tarW  + x;
        int tarArea   = tarW * tarH;

        // bilinear interpolation -- 实现bilinear interpolation + BGR2RGB + shift + nhwc2nchw
        tar[tarIdx + tarArea * 0] = 
            round((a1_1 * src[srcIdx1_1 + 2] + 
                   a1_2 * src[srcIdx1_2 + 2] +
                   a2_1 * src[srcIdx2_1 + 2] +
                   a2_2 * src[srcIdx2_2 + 2])) / 255.0f;

        tar[tarIdx + tarArea * 1] = 
            round((a1_1 * src[srcIdx1_1 + 1] + 
                   a1_2 * src[srcIdx1_2 + 1] +
                   a2_1 * src[srcIdx2_1 + 1] +
                   a2_2 * src[srcIdx2_2 + 1])) / 255.0f;

        tar[tarIdx + tarArea * 2] = 
            round((a1_1 * src[srcIdx1_1 + 0] + 
                   a1_2 * src[srcIdx1_2 + 0] +
                   a2_1 * src[srcIdx2_1 + 0] +
                   a2_2 * src[srcIdx2_2 + 0])) / 255.0f;
    }
}
  1. warpaffine_BGR2RGB_kernel :

  • 执行仿射变换,并将图像从BGR格式转换为RGB格式。

__global__ void warpaffine_BGR2RGB_kernel(
    float* tar, uint8_t* src, 
    TransInfo trans,
    AffineMatrix affine_matrix)
{
    float src_x, src_y;

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    affine_transformation(affine_matrix.reverse, x + 0.5, y + 0.5, &src_x, &src_y);

    int src_x1 = floor(src_x - 0.5);
    int src_y1 = floor(src_y - 0.5);
    int src_x2 = src_x1 + 1;
    int src_y2 = src_y1 + 1;

    if (src_y1 < 0 || src_x1 < 0 || src_y1 > trans.src_h || src_x1 > trans.src_w) {
    } else {
        float tw   = src_x - src_x1;
        float th   = src_y - src_y1;

        float a1_1 = (1.0 - tw) * (1.0 - th);
        float a1_2 = tw * (1.0 - th);
        float a2_1 = (1.0 - tw) * th;
        float a2_2 = tw * th;

        int srcIdx1_1 = (src_y1 * trans.src_w + src_x1) * 3;
        int srcIdx1_2 = (src_y1 * trans.src_w + src_x2) * 3;
        int srcIdx2_1 = (src_y2 * trans.src_w + src_x1) * 3;
        int srcIdx2_2 = (src_y2 * trans.src_w + src_x2) * 3;

        int tarIdx    = y * trans.tar_w  + x;
        int tarArea   = trans.tar_w * trans.tar_h;

        tar[tarIdx + tarArea * 0] = 
            round((a1_1 * src[srcIdx1_1 + 2] + 
                   a1_2 * src[srcIdx1_2 + 2] +
                   a2_1 * src[srcIdx2_1 + 2] +
                   a2_2 * src[srcIdx2_2 + 2])) / 255.0f;

        tar[tarIdx + tarArea * 1] = 
            round((a1_1 * src[srcIdx1_1 + 1] + 
                   a1_2 * src[srcIdx1_2 + 1] +
                   a2_1 * src[srcIdx2_1 + 1] +
                   a2_2 * src[srcIdx2_2 + 1])) / 255.0f;

        tar[tarIdx + tarArea * 2] = 
            round((a1_1 * src[srcIdx1_1 + 0] + 
                   a1_2 * src[srcIdx1_2 + 0] +
                   a2_1 * src[srcIdx2_1 + 0] +
                   a2_2 * src[srcIdx2_2 + 0])) / 255.0f;
    }
}
  1. resize_bilinear_gpu (两个重载版本) :

  • 这两个函数是对CUDA核函数的封装,用于在GPU上执行图像的缩放操作。它们设置CUDA的线程块和网格尺寸,计算缩放比例,并根据指定的策略(如最邻近、双线性插值等)选择相应的核函数进行图像处理。一个版本还包括了像素值归一化的步骤。

这里基本上除了warpffine都是分类网络的前处理, 各种不同的可以拿来做实验

4. 后处理

重点是下面这种写法, 因为vector这种数据结构的类型你在中间删掉了一个,他会动到整个vector, 所以这里bbox会赋予一个flag模式,这样最后是给通过这个查看是否添加进最后的final_bboxes里面, 这样就很高效

vector<bbox> final_bboxes;
    final_bboxes.reserve(m_bboxes.size());
    std::sort(m_bboxes.begin(), m_bboxes.end(), 
              [](bbox& box1, bbox& box2){return box1.confidence > box2.confidence;});

    /*
     * nms在网上有很多实现方法,其中有一些是根据nms的值来动态改变final_bboex的大小(resize, erease)
     * 这里需要注意的是,频繁的对vector的大小的更改的空间复杂度会比较大,所以尽量不要这么做
     * 可以通过给bbox设置skip计算的flg来调整。
    */
    for(int i = 0; i < m_bboxes.size(); i ++){
        if (m_bboxes[i].flg_remove)
            continue;
        
        final_bboxes.emplace_back(m_bboxes[i]);
        for (int j = i + 1; j < m_bboxes.size(); j ++) {
            if (m_bboxes[j].flg_remove)
                continue;

            if (m_bboxes[i].label == m_bboxes[j].label){
                if (iou_calc(m_bboxes[i], m_bboxes[j]) > nms_threshold)
                    m_bboxes[j].flg_remove = true;
            }
        }
    }
    LOGD("the count of bbox after NMS is %d", final_bboxes.size());

5. 后期的提升

  1. 这边的提升思想主要是量化后掉精度的问题, 在这里面韩君给出了一些检查的方案, 因为我们这里是一个Mutil-Task的任务

  • 是否在input/output附近做了int8量化

  • 如果是multi-task的话,是否所有的task都掉点严重

  • calibration的数据集是不是选的不是很好

  • calibration batch size是不是选择的不是很好

  • calibrator是不是没有选择好

  • 某些计算是否不应该做量化

  • 使用polygraphy分析

① 全网独家视频课程

BEV感知 、毫米波雷达视觉融合 多传感器标定 多传感器融合 多模态3D目标检测 点云3D目标检测 目标跟踪 Occupancy、 cuda与TensorRT模型部署 协同感知 语义分割、 自动驾驶仿真、 传感器部署、 决策规划、轨迹预测 等多个方向学习视频( 扫码即可学习

ad73d6172f87a0155934352d335ab0fd.png 视频官网:www.zdjszx.com

② 国内首个自动驾驶学习社区

近2000人的交流社区,涉及30+自动驾驶技术栈学习路线,想要了解更多自动驾驶感知(2D检测、分割、2D/3D车道线、BEV感知、3D目标检测、Occupancy、多传感器融合、多传感器标定、目标跟踪、光流估计)、自动驾驶定位建图(SLAM、高精地图、局部在线地图)、自动驾驶规划控制/轨迹预测等领域技术方案、AI模型部署落地实战、行业动态、岗位发布,欢迎扫描下方二维码,加入自动驾驶之心知识星球, 这是一个真正有干货的地方,与领域大佬交流入门、学习、工作、跳槽上的各类难题,日常分享论文+代码+视频 ,期待交流!

763d7d9d0d935b616abda455e918c111.png

③【自动驾驶之心】技术交流群

自动驾驶之心是首个自动驾驶开发者社区,聚焦 目标检测、语义分割、全景分割、实例分割、关键点检测、车道线、目标跟踪、3D目标检测、BEV感知、多模态感知、Occupancy、多传感器融合、transformer、大模型、点云处理、端到端自动驾驶、SLAM、光流估计、深度估计、轨迹预测、高精地图、NeRF、规划控制、模型部署落地、自动驾驶仿真测试、产品经理、硬件配置、AI求职交流 等方向。扫码添加汽车人助理微信邀请入群,备注:学校/公司+方向+昵称(快速入群方式)

71c909f4e6666db9f00942411eaa8b57.jpeg

④【自动驾驶之心】平台矩阵, 欢迎联系我们!

3d3445399b59363e2d11195fd02550df.jpeg

本文内容由网友自发贡献,版权归原作者所有,本站不承担相应法律责任。如您发现有涉嫌抄袭侵权的内容,请联系:hwhale#tublm.com(使用前将#替换为@)

CUDA与TensorRT | 把检测器加进来,YOLOv8部署实战! 的相关文章

随机推荐

  • QTNet:最新时序融合新方案!点云、图像、多模态检测器全适用(NeurIPS 2023)...

    点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 gt gt 点击进入 自动驾驶之心 多传感器融合 技术交流群 论文作者 Jinghua Hou 编辑 自动驾驶之心 写在前面 个人理解 时序融合能够有效提升自动驾驶3D
  • 免费白嫖GPU算力,畅玩大模型的算力平台!

    要做深度学习 必然需要 GPU 如何构建一个弹性的 GPU 环境是很多读者关心和常问的问题 今天给大家分享 GPU 平台 趋动云 正好官方有活动 新用户 有 价值168元的赠送算力 有效期1年 数量有限先到先得 领取方式见文末 学校课题组没
  • ​NeurIPS 2023|RayDF:实时渲染!基于射线的三维重建新方法

    编辑 极市平台 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 NeRF 技术交流群 本文只做学术分享 如有侵权 联系删文 导读 本论文提出一个全新的维护了多视角几何一致性的基于射线的隐式表达
  • Far3D:直接干到150m,视觉3D目标检测新思路(AAAI2024)

    点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 gt gt 点击进入 自动驾驶之心 3D目标检测 技术交流群 论文作者 自动驾驶Daily 编辑 自动驾驶之心 近来在 Arxiv 读到一篇纯视觉环视感知的新工作 它延
  • 未来之路:大模型技术在自动驾驶的应用与影响

    作者 一颗小树x 编辑 汽车人 原文链接 https zhuanlan zhihu com p 666863252 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 大模型 技术交流群 本文只做
  • 年后跳槽机会多吗?什么时候准备可以实现弯道超车!

    不知不觉2023年即将迎来尾声 秋招也基本上结束了 这几个月汽车人分享了很多自动驾驶和CV方向的工作 像理想 华为的薪资更是羡煞旁人 对于秋招不是很满意的同学和年后打算跳槽的小伙伴 现在可以着手准备起来春招和年后跳槽了 这里也分享一些汽车人
  • 高质量的3D资产如何生成?X-Dreamer或是答案!

    作者 马祎炜 编辑 我爱计算机视觉 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 全栈算法 技术交流群 本文只做学术分享 如有侵权 联系删文 本篇分享论文 X Dreamer Creatin
  • Coco-LIC:紧耦合激光雷达-惯性相机里程计SOTA方案

    点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 gt gt 点击进入 自动驾驶之心 SLAM 技术交流群 论文作者 Xiaolei Lang 编辑 自动驾驶之心 笔者个人理解 机器人在弱结构化 弱纹理环境中的状态估计
  • 超越BEV视角 | 新型紧凑占用Transformer助力自动驾驶3D占用预测

    作者 小书童 编辑 集智书童 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 占用栅格 技术交流群 本文只做学术分享 如有侵权 联系删文 自动驾驶社区对3D占用预测表现出显著兴趣 这主要得益于
  • RV融合!自动驾驶中基于毫米波雷达视觉融合的3D检测综述

    编辑 汽车人 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 多传感器融合 技术交流群 本文只做学术分享 如有侵权 联系删文 自主驾驶在复杂场景下的目标检测任务至关重要 而毫米波雷达和视觉融合
  • 自动驾驶从业者顶流学习笔记,值得收藏!

    自动驾驶之心的兄弟号 自动驾驶Daily后面将会正式投入运营使用 主要关注行业 最新技术分享等 第一次介绍 自动驾驶Daily为大家整理了多个领域方向的数据集 综述 经典论文 视频教程 供大家学习 主要涉及目标检测 语义分割 全景 实例分割
  • 收官,图森即将关闭美国业务

    作者 卡车技术前线 编辑 智车科技 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 本文只做学术分享 如有侵权 联系删文 年末将至 非常遗憾 我们年初的预测成为了现实 自动驾驶卡车创企将面临前所未有的困难 即将进入优
  • 清华大学 | 智能产业研究院自动驾驶冬令营学员招募!

    点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 gt gt 点击进入 自动驾驶之心 求职交流 技术交流群 论文作者 汽车人 编辑 自动驾驶之心 AIR 自动驾驶方向冬令营 在这个冬季 清华大学智能产业研究院 AIR
  • 硬核 | 从零制作一个激光雷达需要多久?

    编辑 ADS智库 点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 点击进入 自动驾驶之心 硬件交流 技术交流群 本文只做学术分享 如有侵权 联系删文 激光雷达 LiDAR 是激光探测及测距系统的简称 目前广泛应用在无
  • Solidity之旅(七)单位以及全局变量

    01 以太币单位 Ether 甭管是虚拟货币还是现实中的 稳定币 他们都有各自的货币单位 当然咯 作为以太坊区块链上交易的虚拟货币以太坊也是有属于自己的那一套货币单位的 而目前以太币 Ether 主要分为这三个 wei gwei以及ethe
  • 如何对LoRA模型的多址接入进行技术研究与改进

    LoRA技术的多址接入是物联网通信中的重要环节 通过对多址接入技术的研究与改进 可以提高LoRA模型的通信容量和效率 本文将详细介绍LoRA技术的多址接入原理 以及当前存在的问题 并提出一些可能的技术研究和改进方向 一 LoRA技术的多址接
  • 【ZYNQ学习】PL第一课

    这节课讲什么 这节课的名字本来是想写为LED 但这一课里除了LED也有按键 又想换为GPIO控制 但关于PL的GPIO控制 不应该这么草率和简单 而且这一课有很多和ZYNQ或者PL关联性不强的东西要说 所以我写了删删了写改了好几遍 终于定为
  • 推荐几个计算机视觉与自动驾驶相关的平台!

    今年来 各家自动驾驶与AI公司开始规模化量产 可落地的技术成为大家争先占领的重点 然而这个行业对从业者能力要求较高 内部非常卷 一个岗位难求 如何从内卷中脱颖而出 除了极强的自律外 系统的学习方法也很重要 这里给大家推荐了几个国内非常具有影
  • Linux进程通信之管道:实现进程间数据传输和共享

    管道是Linux进程间通信的一种重要方式 它通过pipe inode文件节点来实现 这个文件节点基于pipefs文件系统 它是一个伪文件系统 专门用于管道的实现 为了更好地理解管道的工作原理 我们可以将pipe inode节点分为三个部分
  • CUDA与TensorRT | 把检测器加进来,YOLOv8部署实战!

    点击下方 卡片 关注 自动驾驶之心 公众号 ADAS巨卷干货 即可获取 gt gt 点击进入 自动驾驶之心 模型部署 技术交流群 论文作者 汽车人 编辑 自动驾驶之心 0 把检测器加进来 本文是我在学习韩博 CUDA与TensorRT部署实