七. 部署YOLOv8检测器-affine-transformation

前言

自动驾驶之心推出的 《CUDA与TensorRT部署实战课程》，链接。记录下个人学习笔记，仅供自己参考

本次课程我们来学习课程第七章—部署YOLOv8检测器，一起来学习仿射变换

课程大纲可以看下面的思维导图

0. 简述

本小节目标：学习 affine transformation 仿射变换

这节我们主要学习仿射变换 warpAffine，理解仿射变换矩阵及其逆矩阵

下面我们开始本次课程的学习🤗

1. 案例运行

在正式开始课程之前，博主先带大家跑通 7.2-affine-transformation 这个小节的案例🤗

源代码获取地址：https://github.com/kalfazed/tensorrt_starter

首先大家需要把 tensorrt_starter 这个项目给 clone 下来，指令如下：

git clone https://github.com/kalfazed/tensorrt_starter.git

也可手动点击下载，点击右上角的 Code 按键，将代码下载下来。至此整个项目就已经准备好了。也可以点击 here 下载博主准备好的源代码（注意代码下载于 2024/7/14 日，若有改动请参考最新）

整个项目后续需要使用的软件主要有 CUDA、cuDNN、TensorRT、OpenCV，大家可以参考 Ubuntu20.04软件安装大全 进行相应软件的安装，博主这里不再赘述

假设你的项目、环境准备完成，下面我们一起来运行下 7.2-affine-transformation 小节案例代码

开始之前我们需要在 tensorrt_starter/chapter7-deploy-yolo-detection/7.2-affine-transformation 小节中创建一个 results 文件夹用于保存仿射变换后的图片

创建完后 7.2 小节整个目录结构如下：

接着我们需要运行代码来执行仿射变换，在此之前我们需要修改下整体的 Makefile.config，指定一些库的路径：

# tensorrt_starter/config/Makefile.config
# CUDA_VER                    :=  11
CUDA_VER                    :=  11.6
    
# opencv和TensorRT的安装目录
OPENCV_INSTALL_DIR          :=  /usr/local/include/opencv4
# TENSORRT_INSTALL_DIR        :=  /mnt/packages/TensorRT-8.4.1.5
TENSORRT_INSTALL_DIR        :=  /home/jarvis/lean/TensorRT-8.6.1.6

Note：大家查看自己的 CUDA 是多少版本，修改为对应版本即可，另外 OpenCV 和 TensorRT 修改为你自己安装的路径即可

接着我们就可以来执行编译，指令如下：

make -j64

输出如下：

接着执行：

./bin/trt-cuda

输出如下：

在 results 文件夹下我们可以看到经过各个预处理方式后的图片，大家可以查看，如下图所示：

如果大家能够看到上述输出结果，那就说明本小节案例已经跑通，下面我们就来看看具体的代码实现

2. 补充说明

在分析代码之前我们先来看下韩君老师在这小节中写的 README 文档

这个小节是接着第二章节的 2.10-bilinear-interpolation 的扩展，大家感兴趣的话可以看看：二. CUDA编程入门-双线性插值计算

之前在 classification model 推理的时候，我们只需要对图像做一次 resize 就好了，但是在做 detection model 例如 yolo 推理的时候，我们首先需要将图片 resize 到 yolo model 可以识别的大小（例如 640x640），之后我们得到在这个尺寸下的 bbox，但是我们在绘图的时候还需要将 bbox 还原成原图的大小，所以需要某一种形式进行 resize 的正向/反向的变换，这个可以通过 warpAffine 仿射变换来实现

其实我们前面 bilinear-interpolation 案例的实现已经非常贴近 warpAffine 了，部分代码如下所示：

// 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;

...

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

warpAffine 的基本公式如下：

// forward
forward_scale = min(tar_w / src_w, tar_h / src_h);
tar_x         = src_x * forward_scale + forward_shift_x;
tar_y         = src_y * forward_scale + forward_shift_y;

// reverse
reverse_scale   = 1 / forward_scale;
reverse_shift_x = -forward_shift / forward_scale_x;
reverse_shift_y = -forward_shift / forward_scale_y;
src_x           = tar_x * reverse_scale + reverse_shift_x;
src_y           = tar_y * reverse_scale + reverse_shift_y;

规范之后可以写成下面这种形式：

// forward
tar_x  =  src_x * forward_scale + src_y * 0             + forward_shift_x;
tar_y  =  src_x * 0             + src_y * forward_scale + forward_shift_y;

// reverse
src_x  =  tar_x * reverse_scale + tar_y * 0             + reverse_shift_x;
src_y  =  tar_x * 0             + tar_y * reverse_scale + reverse_shift_y;

我们可以通过 Matrix 的形式保存这些 scale 和 shift 等需要的时候直接使用，在 yolo 的 preprocess 中由于我们需要把图片 resize 成 letterbox 所以可以把 scale 和 shift 写成下面的形式：

// 存储forward时需要的scale和shift
void calc_forward_matrix(TransInfo trans){
    forward[0] = forward_scale;
    forward[1] = 0;
    forward[2] = - forward_scale * trans.src_w * 0.5 + trans.tar_w * 0.5;
    forward[3] = 0;
    forward[4] = forward_scale;
    forward[5] = - forward_scale * trans.src_h * 0.5 + trans.tar_h * 0.5;
};

// 存储reverse时需要的scale和shift
void calc_reverse_matrix(TransInfo trans){
    reverse[0] = reverse_scale;
    reverse[1] = 0;
    reverse[2] = - reverse_scale * trans.tar_w * 0.5 + trans.src_w * 0.5;
    reverse[3] = 0;
    reverse[4] = reverse_scale;
    reverse[5] = - reverse_scale * trans.tar_h * 0.5 + trans.src_h * 0.5;
};

// 仿射变换的计算公式
__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];
}

Note：关于仿射变换的更多细节大家感兴趣的可以看看杜老师的讲解视频以及博主之前写的文章：

• 手写AI/预处理高性能
• YOLOv5推理详解及预处理高性能实现

大家可能看上面的 warpAffine 的代码有些困惑，那其实对照着公式来看还是比较清晰的，博主这边还是快速过一遍吧

在目标检测任务中，我们的预处理操作通常是先对图像进行等比缩放，然后居中，多余部分填充，类似于下图所展示的，这个过程也被叫做 letterbox 添加灰条：

整个过程可以拆分为以下三个步骤：

1. 等比缩放，矩阵 $S$ 实现（缩放倍数 scale 为目标图像与源图像宽比值和高比值的最小值）

2. 将图片中心平移到左上角坐标原点，矩阵 $O$ 实现

3. 将图片平移到目标位置的中心，矩阵 $T$ 实现

预处理的过程可以通过三个矩阵完成，分别为缩放矩阵 $S$，平移矩阵 $O$，平移矩阵 $T$，将这三个矩阵可以进行合并成一个矩阵 $M(M=TOS)$，其中 $M$ 矩阵被称为仿射变换矩阵，可以帮助我们完成图像预处理工作

我们可以直接写出仿射变换矩阵 $M$ 的计算公式，如下式 (4) 所示，其中 $Origin$ 代表源图像，$Dst$ 代表目标图像，$(x,y)$ 为源图像上任意像素点坐标，$(x^{\prime },y^{\prime })$ 为目标图像上任意像素点坐标

通过上述变换可以得到源图像任意像素值的坐标对应在目标图像的位置，那么整个预处理过程也就可以通过仿射变换矩阵 $M$ 来完成。除了将源图像通过 $M$ 变换到目标图像，我们也要考虑将目标图像映射回源图像，即求 $M$ 矩阵的逆矩阵 $M^{-1}$

逆变换矩阵 $M^{-1}$ 计算公式如下：

设

则有

OK，以上就是 warpAffine Matrix 的推导公式，理解了之后大家再来看之前的代码会发现还是比较容易简单的：

// forward
tar_x  =  src_x * forward_scale + src_y * 0             + forward_shift_x;
tar_y  =  src_x * 0             + src_y * forward_scale + forward_shift_y;

// reverse
src_x  =  tar_x * reverse_scale + tar_y * 0             + reverse_shift_x;
src_y  =  tar_x * 0             + tar_y * reverse_scale + reverse_shift_y;

理解了上面的内容我们下面来看代码就轻松很多了

3. 代码分析

3.1 main.cpp

我们先从 main.cpp 看起：

#include <stdio.h>
#include <cuda_runtime.h>
#include <iostream>

#include "utils.hpp"
#include "timer.hpp"
#include "preprocess.hpp"

using namespace std;

int main(){
    Timer timer;

    string file_path     = "data/deer.png";
    string output_prefix = "results/";
    string output_path   = "";

    cv::Mat input = cv::imread(file_path);
    int tar_h = 500;
    int tar_w = 550;
    int tactis;

    cv::Mat resizedInput_cpu;
    cv::Mat resizedInput_gpu;
    
    /* 
     * bilinear interpolation resize的CPU/GPU速度比较
     * 由于CPU端做完预处理之后，进入如果有DNN也需要将数据传送到device上，
     * 所以这里为了让测速公平，仅对下面的部分进行测速:
     *
     * - host端
     *     cv::resize的bilinear interpolation
     *     normalization进行归一化处理
     *     BGR2RGB来实现通道调换
     *
     * - device端
     *     bilinear interpolation + normalization + BGR2RGB的自定义核函数
     *
     * 由于这个章节仅是初步CUDA学习，大家在自己构建推理模型的时候可以将这些地方进行封装来写的好看点，
     * 在这里代码我们更关注实现的逻辑部分
     *
     * tatics 列表
     * 0: 最近邻差值缩放 + 全图填充
     * 1: 双线性差值缩放 + 全图填充
     * 2: 双线性差值缩放 + 填充(letter box)
     * 3: 双线性差值缩放 + 填充(letter box) + 平移居中
     * 4: 仿射变换(letter box)
     * */
    
    resizedInput_cpu = preprocess_cpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_bilinear_cpu.png";
    cv::cvtColor(resizedInput_cpu, resizedInput_cpu, cv::COLOR_RGB2BGR);
    cv::imwrite(output_path, resizedInput_cpu);

    tactis = 0;
    resizedInput_gpu = preprocess_gpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_nearest_gpu.png";
    cv::imwrite(output_path, resizedInput_gpu);

    tactis = 1;
    resizedInput_gpu = preprocess_gpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_bilinear_gpu.png";
    cv::imwrite(output_path, resizedInput_gpu);

    tactis = 2;
    resizedInput_gpu = preprocess_gpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_bilinear_letterbox_gpu.png";
    cv::imwrite(output_path, resizedInput_gpu);

    tactis = 3;
    resizedInput_gpu = preprocess_gpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_bilinear_letterbox_center_gpu.png";
    cv::imwrite(output_path, resizedInput_gpu);

    tactis = 4;
    resizedInput_gpu = preprocess_gpu(input, tar_h, tar_w, timer, tactis);
    output_path = output_prefix + getPrefix(file_path) + "_resized_warpaffine_letterbox_center_gpu.png";
    cv::imwrite(output_path, resizedInput_gpu);

    return 0;
}

在 main 函数中我们对输入图像进行了各种预处理操作，并比较了 CPU 和 GPU 端的图像缩放、归一化、通道交换等操作的性能差异，最后将每个策略的处理结果保存为不同的文件

3.2 preprocess.cu

在 main 函数中的 preprocess_gpu 会通过不同的 tactics 调用不同的 Kernel 核函数完成预处理操作，主要包括以下四种：

void resize_bilinear_gpu(
    uint8_t* d_tar, uint8_t* d_src, 
    int tarW, int tarH, 
    int srcW, int srcH, 
    int tactis) 
{
    dim3 dimBlock(16, 16, 1);
    dim3 dimGrid(tarW / 16 + 1, tarH / 16 + 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);

    if (tactis > 1) {
        scaled_h = scale;
        scaled_w = scale;
    }

    // for affine transformation
    TransInfo    trans(srcW, srcH, tarW, tarH);
    AffineMatrix affine;
    affine.init(trans);
    
    switch (tactis) {
    case 0:
        resize_nearest_BGR2RGB_kernel <<<dimGrid, dimBlock>>> (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case 1:
        resize_bilinear_BGR2RGB_kernel <<<dimGrid, dimBlock>>> (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case 2:
        resize_bilinear_BGR2RGB_kernel <<<dimGrid, dimBlock>>> (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case 3:
        resize_bilinear_BGR2RGB_shift_kernel <<<dimGrid, dimBlock>>> (d_tar, d_src, tarW, tarH, srcW, srcH, scaled_w, scaled_h);
        break;
    case 4:
        resize_warpaffine_BGR2RGB_kernel <<<dimGrid, dimBlock>>> (d_tar, d_src, trans, affine);
        break;
    default:
        break;
    }
}

前面三种策略我们在 2.10-bilinear-interpolation 小节案例中已经详细分析过了，这边博主就不再赘述了，我们重点来看第四种策略，也就是 warpAffine 仿射变换的 Kernel 核函数实现，完整的代码如下所示：

struct TransInfo{
    int src_w;
    int src_h;
    int tar_w;
    int tar_h;
    TransInfo(int srcW, int srcH, int tarW, int tarH):
        src_w(srcW), src_h(srcH), tar_w(tarW), tar_h(tarH){}
};

struct AffineMatrix{
    float forward[6];
    float reverse[6];
    float forward_scale;
    float reverse_scale;

    void calc_forward_matrix(TransInfo trans){
        forward[0] = forward_scale;
        forward[1] = 0;
        forward[2] = - forward_scale * trans.src_w * 0.5 + trans.tar_w * 0.5;
        forward[3] = 0;
        forward[4] = forward_scale;
        forward[5] = - forward_scale * trans.src_h * 0.5 + trans.tar_h * 0.5;
    };

    void calc_reverse_matrix(TransInfo trans){
        reverse[0] = reverse_scale;
        reverse[1] = 0;
        reverse[2] = - reverse_scale * trans.tar_w * 0.5 + trans.src_w * 0.5;
        reverse[3] = 0;
        reverse[4] = reverse_scale;
        reverse[5] = - reverse_scale * trans.tar_h * 0.5 + trans.src_h * 0.5;
    };

    void init(TransInfo trans){
        float scaled_w = (float)trans.tar_w / trans.src_w;
        float scaled_h = (float)trans.tar_h / trans.src_h;
        forward_scale = (scaled_w < scaled_h ? scaled_w : scaled_h);
        reverse_scale = 1 / forward_scale;
    
        // 计算src->tar和tar->src的仿射矩阵
        calc_forward_matrix(trans);
        calc_reverse_matrix(trans);
    }
};

__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];
}

__global__ void resize_warpaffine_BGR2RGB_kernel(
    uint8_t*     tar, 
    uint8_t*     src, 
    TransInfo    trans_info,
    AffineMatrix matrix)
{
    float src_x, src_y;

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

    // bilinear interpolation -- 通过逆仿射变换得到计算tar中的x, y所需要的src中的src_x, src_y
    affine_transformation(matrix.reverse, x + 0.5, y + 0.5, &src_x, &src_y);

    // bilinear interpolation -- 计算x,y映射到原图时最近的4个坐标
    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_info.src_h || src_x1 > trans_info.src_w) {
        // bilinear interpolation -- 对于越界的坐标不进行计算
    } else {
        // bilinear interpolation -- 计算原图上的坐标(浮点类型)在0~1之间的值
        float tw   = src_x - src_x1;
        float th   = src_y - src_y1;

        // 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 * trans_info.src_w + src_x1) * 3;  //左上
        int srcIdx1_2 = (src_y1 * trans_info.src_w + src_x2) * 3;  //右上
        int srcIdx2_1 = (src_y2 * trans_info.src_w + src_x1) * 3;  //左下
        int srcIdx2_2 = (src_y2 * trans_info.src_w + src_x2) * 3;  //右下

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


        // bilinear interpolation -- 实现bilinear interpolation + BGR2RGB
        tar[tarIdx + 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]);

        tar[tarIdx + 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]);

        tar[tarIdx + 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]);
    }
}

这个 CUDA 核函数实现了图像的 warpAffine 仿射变换操作，结合 bilinear interpolation 双线性插值来对图像进行缩放，并同时执行 BGR 到 RGB 的颜色通道转换，下面我们来详细分析下这个核函数的各个部分：(from ChatGPT)

首先我们来看核函数的各个参数：

__global__ void resize_warpaffine_BGR2RGB_kernel(
    uint8_t*     tar, 
    uint8_t*     src, 
    TransInfo    trans_info,
    AffineMatrix matrix)

• tar：目标图像的存储数组，输出图像数据会保存在这里
• src：源图像的存储数组，即输入图像的数据
• trans_info：包含源图像和目标图像的宽度和高度的结构体
• matrix：包含仿射变换矩阵以及正向和反向缩放系数的结构体

TransInfo 结构体定义如下：

struct TransInfo{
    int src_w;
    int src_h;
    int tar_w;
    int tar_h;
    TransInfo(int srcW, int srcH, int tarW, int tarH):
        src_w(srcW), src_h(srcH), tar_w(tarW), tar_h(tarH){}
};

功能：

• TransInfo 结构体用于存储源图像和目标图像的宽度和高度信息
• 它有四个成员变量：
  • src_w：源图像的宽度
  • src_h：源图像的高度
  • tar_w：目标图像的宽度
  • tar_h：目标图像的高度
• 通过构造函数 TransInfo(int srcW, int srcH, int tarW, int tarH) 可以快速初始化这些参数

用途：

• 该结构体主要用于传递源图像和目标图像的尺寸信息，以便于计算仿射变换矩阵和进行图像缩放

AffineMatrix 结构体定义如下：

struct AffineMatrix{
    float forward[6];
    float reverse[6];
    float forward_scale;
    float reverse_scale;

    void calc_forward_matrix(TransInfo trans){
        forward[0] = forward_scale;
        forward[1] = 0;
        forward[2] = - forward_scale * trans.src_w * 0.5 + trans.tar_w * 0.5;
        forward[3] = 0;
        forward[4] = forward_scale;
        forward[5] = - forward_scale * trans.src_h * 0.5 + trans.tar_h * 0.5;
    };

    void calc_reverse_matrix(TransInfo trans){
        reverse[0] = reverse_scale;
        reverse[1] = 0;
        reverse[2] = - reverse_scale * trans.tar_w * 0.5 + trans.src_w * 0.5;
        reverse[3] = 0;
        reverse[4] = reverse_scale;
        reverse[5] = - reverse_scale * trans.tar_h * 0.5 + trans.src_h * 0.5;
    };

    void init(TransInfo trans){
        float scaled_w = (float)trans.tar_w / trans.src_w;
        float scaled_h = (float)trans.tar_h / trans.src_h;
        forward_scale = (scaled_w < scaled_h ? scaled_w : scaled_h);
        reverse_scale = 1 / forward_scale;
    
        // 计算src->tar和tar->src的仿射矩阵
        calc_forward_matrix(trans);
        calc_reverse_matrix(trans);
    }
};

功能：

• AffineMatrix 结构体用于计算和存储源图像和目标图像之间的仿射变换矩阵及其逆矩阵，以及对应的缩放因子
• 主要成员变量：
  • foward[6]：正向仿射变换矩阵（源图像到目标图像）
  • reverse[6]：反向仿射变换矩阵（目标图像到源图像）
  • forward_scale：正向缩放系数
  • reverse_scale：反向缩放系数

主要方法：

• calc_forward_matrix(TransInfo trans)：
  • 计算正向仿射变换矩阵，即从源图像到目标图像的变换
  • 使用 forward_scale 来调整源图像的大小，并通过调整平移项确保图像居中
• calc_reverse_matrix(TransInfo trans)：
  • 计算反向仿射变换矩阵，即从目标图像到源图像的变换。该矩阵用于将目标图像中的像素位置映射回源图像，以便插值计算
  • 通过反向缩放系数调整目标图像的大小，使其匹配源图像
• init(TransInfo trans)：
  • 初始化正向和反向的缩放比例，并调用 calc_forward_matrix 和 calc_reverse_matrix 计算正向和反向的仿射变换矩阵
  • forward_scale 是根据目标图像和源图像的宽度和高度的比例来计算的，选择其中较小的缩放因子，这样可以保持纵横比
  • reverse_scale 是 forward_scale 的倒数，保证反向映射的正确性

用途：

• AffineMatrix 是图像缩放和仿射变换的核心，它的正向矩阵用于图像的放缩操作，反向矩阵用于逆仿射变换（即将目标图像中的坐标映射到源图像中）
• 通过这个矩阵，我们可以完成图像缩放、旋转、平移等操作

接着我们来看核函数的工作流程，核函数的主要流程是：

• 使用逆仿射变换计算目标图像像素在源图像中的对应位置
• 使用双线性插值计算目标图像的像素值，基于源图像四个相邻像素点的加权平均
• 将像素值从 BGR 格式转换为 RGB 格式

1. 计算目标图像的坐标 x 和 y

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

• 这部分代码利用 CUDA 的线程模型，通过 blockIdx 和 threadIdx 来计算每个线程所处理的目标图像中的像素坐标 (x, y)

2. 逆仿射变换计算源图像的坐标

• 调用 affine_transformation 函数，利用逆仿射变换矩阵 matrix.reverse 将目标图像坐标 (x, y) 转换为源图像中的浮点坐标 (src_x, src_y)

affine_transformation 函数的定义如下：

__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];
}

功能：

• affine_transformation 函数用于在 device 上执行仿射变换，它将源图像中的 (src_x, src_y) 坐标转换为目标图像中的 (tar_x, tar_y) 坐标
• 该函数通过乘以 2x3 的仿射变换矩阵来完成坐标转换：
  • tar_x = M[0] * src_x + M[1] * src_y + M[2]
  • tar_y = M[3] * src_x + M[4] * src_y + M[5]

详细说明：

• 仿射变换矩阵 trans_matrix[6] 是一个 2x3 的矩阵，它定义了如何将源图像中的坐标映射到目标图像中的坐标，具体公式如下：
  • $x^{\prime }=M_{00}\cdot x+M_{01}\cdot y+M_{02}$
  • $y^{\prime }=M_{10}\cdot x+M_{11}\cdot y+M_{12}$
• trans_matrix[6] 包含了以下值：
  • M[0] 对应 forward[0] 或 reverse[0]，用于缩放
  • M[1] 和 M[3] 用于旋转
  • M[2] 和 M[5] 是平移项，确保图像的中心对齐

用途：

• 该函数在核函数中被调用，通过传入反向仿射矩阵 matrix.reverse，将目标图像的像素坐标映射到源图像中，这是进行双线性插值的基础操作

3. 计算源图像最近的四个相邻像素

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;

• 使用双线性插值，需要计算源图像中与浮点坐标 (src_x, src_y) 最接近的四个整数坐标像素，分别为 (src_x1, src_y1)、(src_x2, src_y1)、(src_x1, src_y2)、(src_x2, src_y2)

4. 边界检查

if (src_y1 < 0 || src_x1 < 0 || src_y1 > trans_info.src_h || src_x1 > trans_info.src_w) {
    // bilinear interpolation -- 对于越界的坐标不进行计算
}

• 这段代码检查源图像中的四个相邻像素点是否越界（即坐标超出了图像的范围），如果越界，则跳过当前像素的计算

5. 双线性插值权重计算

float tw   = src_x - src_x1;
float th   = src_y - src_y1;

• 计算 src_x 和 src_y 与它们左上角元素 (src_x1, src_y1) 的距离，分别为 tw 和 th，这将用于计算插值权重

6. 计算每个邻近像素的插值权重

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;                  //左上

• 通过双线性插值计算四个相邻像素的权重，插值权重的计算公式基于源图像中相邻像素与目标像素的距离
  • a1_1：左上角像素的权重
  • a1_2：右上角像素的权重
  • a2_1：左下角像素的权重
  • a2_2：右下角像素的权重

7. 计算源图像中四个像素点的索引

int srcIdx1_1 = (src_y1 * trans_info.src_w + src_x1) * 3;  //左上
int srcIdx1_2 = (src_y1 * trans_info.src_w + src_x2) * 3;  //右上
int srcIdx2_1 = (src_y2 * trans_info.src_w + src_x1) * 3;  //左下
int srcIdx2_2 = (src_y2 * trans_info.src_w + src_x2) * 3;  //右下

• 计算源图像四个相邻像素点在 src 数组中的索引，这里每个像素点有 RGB 三个通道，所以每个索引都乘以 3

8. 计算目标图像中当前像素的索引

int tarIdx = (y * trans_info.tar_w + x) * 3;

• 计算目标图像中当前像素在 tar 数组中的索引，同样，每个像素有 RGB 三个通道，所以索引乘以 3

9. 双线性插值计算 + BGR2RGB

tar[tarIdx + 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]);

tar[tarIdx + 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]);

tar[tarIdx + 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]);

• 利用前面计算出的权重和源图像四个像素点的值，对每个颜色通道（B、G、R）进行双线性插值计算，值得注意的是在这个过程中，通道顺序从 BGR 转换为了 RGB
  • tar[tarIdx + 0]：R 通道
  • tar[tarIdx + 1]：G 通道
  • tar[tarIdx + 2]：B 通道

总的来说，这个 CUDA 核函数实现了基于仿射变换的双线性插值，并且将源图像的 BGR 颜色通道转换为目标图像的 RGB 颜色通道

OK，以上就是 warpAffine 实现代码的详细分析了，理解原理之后再来看代码还是比较简单的

结语

本次课程我们学习了 affine transformation 仿射变换，并简单分析了其具体实现的代码，warpAffine 在许多模型的前处理中都有使用到，因此理解它是必不可少的

OK，以上就是 7.2 小节案例的全部内容了，下节我们来学习 7.3 小节 yolov8 模型的部署，敬请期待😄

下载链接

• tensorrt_starter源码

参考

• Ubuntu20.04软件安装大全
• https://github.com/kalfazed/tensorrt_starter
• 二. CUDA编程入门-双线性插值计算
• 手写AI/预处理高性能
• YOLOv5推理详解及预处理高性能实现