前言

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

Note：关于 CUDA 加速双线程插值的内容博主之前有简单记录过，感兴趣的可以看看 YOLOv5推理详解及预处理高性能实现

本次课程我们来学习课程第二章—CUDA 编程入门，一起来学习双线性插值的计算

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

0. 简述

本小节目标：理解如何使用 cuda 进行 opencv 的图像处理的加速，理解双线性插值进行图像大小调整的算法流程

这节我们来讲第二章第 5 小节，双线性插值的计算，这个小节的案例更偏实际应用，大家在利用 TensorRT 做模型部署的时候会考虑将输入的图片统一缩放到同一尺寸大小，如果我们直接使用 OpenCV 的 resize 函数其实效率并不高，那我们这个小节就要学习如何利用 CUDA 的高并发行特性对双线性插值进行加速

1. 执行一下我们的第十个CUDA程序

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

这节课的案例代码是 2.10-bilinear-interpolation，如下所示：

这个小节的案例代码核心是 preprocess.cpp 和 preprocess.cu 两个文件，preprocess.cpp 中是 bilinear resize 和 opencv 实现以及 bilinear resize 的 CUDA 接口实现，如下图所示：

preprocess.cpp

而 preprocess.cu 则是利用 CUDA 核函数来实现双线程插值，如下所示：

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

    // 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_y1 > srcH || src_x1 > 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) * 3;

        // bilinear interpolation -- 实现bilinear interpolation的resize + 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]);
    }
}

preprocess.cu 中使用了多个核函数来实现双线性插值，插值效果对比如下图所示：

resized_bilinear_gpu 是一个普通的双线性插值，它将图片缩放到对应的尺寸且不做填充，resized_bilinear_letterbox_gpu 则是保证缩放后图片的长宽比保持一致多余部分填充，resized_bilinear_letterbox_center_gpu 则是将 letterbox 后的图片进行居中，一般来说我们在分类模型的预处理中使用 resized_bilinear_gpu 较多，而在检测模型的预处理中使用 resized_bilinear_letterbox_center_gpu 较多

本节案例执行效果如下：

我们可以看到不同双线性插值实现的时间对比，一共有 5 个数据，第一个是 CPU 上双线性插值的执行时间大概是 1.18ms，其次是 GPU 上最近邻插值的执行时间大概是 0.021ms，接着是 GPU 上双线性插值的时间大概是 0.037ms，最后两个是双线性插值-letterbox 的执行时间大概是 0.018ms，相比于 CPU 上的执行时间大概可以加速 60 倍，这个还是比较夸张的，大家可以自己更改 kernel 核函数的一些参数看下实际效果

2. Bilinear interpolation

在代码分析之前我们有必要聊下 Bilinear interpolation 双线性插值，那双线性插值到底是什么呢？它其实是一种对图像进行缩小/放大的计算方法，也是 opencv 默认的 resize 方式

以上图为例来讲解，现在我们这里有一张 1000x600 分辨率的狐狸图片，我们需要将它缩小到 256x256 分辨率，我们怎么来实现呢？那其实这里是有一个映射关系存在的（具体推导大家可以参考 here），也就是说对于目标图像 dst 上任意像素点我们都可以通过映射关系找到它在源图像 src 上的位置，从而将源图像上的像素点填充到目标图像上，如下图所示：

但是这里存在一个问题，那就是目标图像通过映射关系到源图像上时的位置坐标不一定是整数，这也就是意味着我们没有办法直接获取到对应位置的像素值，这个时候我们就可以利用插值算法来取像素，如下图所示：

现在假设我们目标图像 dst 的 (100,100) 处要填充的像素值映射回源图像 src 时的位置是 (200.14,100.11)，由于映射回源图时的位置坐标是浮点数导致我们没有办法直接获取像素值，但是我们可以得到距离它最近的周围四个点的像素值即 (200,100)、(200,101)、(201,100) 以及 (201,101)

它们四个点的像素值是已知的，那中间红色点 (200.14,100.11) 像素值该怎么求呢？我们可以通过加权平均的方式来求取即

$$dst[100,100]=src[100,200]\times a11+src[101,200]\times a10+src[100,201]\times a01+src[101,201]\times a00$$

这个其实就是双线性插值的过程，通过未知像素点周围的四个像素来进行加权求和得到最终的像素，那依此类推最近邻插值是怎么做的呢？那通过名字我们就可以知道这个插值方式其实就是选取距离未知像素点最近的那个像素点的像素值，最近邻插值计算速度相较于双线性插值要快但图像质量相较于双线性插值要差

如果我们想要缩放后的图像仍然保持相同的长宽比，我们需要保证图像的宽和高的缩放比一致，最终的缩放因子 $scale=min(\frac{tar\_w}{src\_w},\frac{tar\_h}{src\_h})$ 最终实现的效果如下图所示：

一般来说在进行缩放以后我们还希望图像居中，所以还需要让图像的中心坐标 shift 一定的像素，如下图所示：

shift 代码如下：

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

在代码中我们让图像的中心点坐标平移到最上面之后再移动到整个目标图像的中心，那其实以上的各种缩放、平移变换我们其实可以直接通过一个仿射变换矩阵解决，这个我们在 YOLOv5推理详解及预处理高性能实现 中有提到过，大家感兴趣的可以看看

利用 CUDA 来实现双线性插值其实还有其它的好处，由于 CUDA 中我们是启动多个线程处理一张图像，每个线程处理一个像素是像素级别的处理，因此很容易实现 BGR2RGB、减均值除标准差等操作，也就是说我们可以通过一个核函数把图像预处理操作全都给做了，这个我们在之后 affine_transformation（仿射变换）的案例再讲解

3. 代码分析

OK，下面我们一起来看下代码，先从 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 = 250;
    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) + 平移居中
     * */
    
    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::cvtColor(resizedInput_cpu, resizedInput_cpu, cv::COLOR_RGB2BGR);
    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::cvtColor(resizedInput_cpu, resizedInput_cpu, cv::COLOR_RGB2BGR);
    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::cvtColor(resizedInput_cpu, resizedInput_cpu, cv::COLOR_RGB2BGR);
    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::cvtColor(resizedInput_cpu, resizedInput_cpu, cv::COLOR_RGB2BGR);
    cv::imwrite(output_path, resizedInput_gpu);
    return 0;
}

这段代码主要实现了一个用于图像处理的程序，旨在比较CPU和GPU在进行双线性插值图像缩放时的性能。代码通过分别使用CPU和GPU来处理同一图像，并将处理结果保存下来，以便于性能比较。下面是对代码的简单分析：(from chatGPT)

主函数（main）定义：

• 初始化一个 Timer 类的实例，用于测量代码执行时间。
• 定义了几个字符串变量来指定输入文件的路径、输出文件的前缀和完整输出路径。
• 读取指定路径的图像文件到 cv::Mat 类型的变量 input 中，这里使用的是 OpenCV 库的imread函数。
• 定义了目标高度 tar_h 和目标宽度 tar_w 变量，以及 tactis 变量用于指定缩放策略。

图像处理逻辑：

• 对输入图像使用 CPU 进行预处理，包括调整大小、归一化和通道转换。处理后的图像保存在 resizedInput_cpu 中。
• 使用 cv::cvtColor 函数将处理后的图像从 RGB 转换为 BGR 格式，这通常是因为 OpenCV 默认使用 BGR 颜色空间。
• 使用 cv::imwrite 函数将处理后的图像写入到文件系统中。
• 通过变化 tactis 的值，分别使用 GPU 进行不同的预处理操作，并保存不同的输出文件。预处理包括最近邻缩放、双线性缩放、带填充的双线性缩放等。

接着我们来看下 preprocess.cpp，代码如下：

#include "preprocess.hpp"
#include "opencv2/opencv.hpp"
#include "utils.hpp"
#include "timer.hpp"

// 根据比例进行缩放 (CPU版本)
cv::Mat preprocess_cpu(cv::Mat &src, const int &tar_h, const int &tar_w, Timer timer, int tactis) {
    cv::Mat tar;

    int height  = src.rows;
    int width   = src.cols;
    float dim   = std::max(height, width);
    int resizeH = ((height / dim) * tar_h);
    int resizeW = ((width / dim) * tar_w);

    int xOffSet = (tar_w - resizeW) / 2;
    int yOffSet = (tar_h - resizeH) / 2;

    resizeW    = tar_w;
    resizeH    = tar_h;

    timer.start_cpu();

    /*BGR2RGB*/
    cv::cvtColor(src, src, cv::COLOR_BGR2RGB);

    /*Resize*/
    cv::resize(src, tar, cv::Size(resizeW, resizeH), 0, 0, cv::INTER_LINEAR);

    timer.stop_cpu();
    timer.duration_cpu<Timer::ms>("Resize(bilinear) in cpu takes:");

    return tar;
}

// 根据比例进行缩放 (GPU版本)
cv::Mat preprocess_gpu(
    cv::Mat &h_src, const int& tar_h, const int& tar_w, Timer timer, int tactis) 
{
    uint8_t* d_tar = nullptr;
    uint8_t* d_src = nullptr;

    cv::Mat h_tar(cv::Size(tar_w, tar_h), CV_8UC3);

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

    int src_size  = height * width * chan;
    int tar_size  = tar_h * tar_w * chan;

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

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

    timer.start_gpu();

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

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

    timer.stop_gpu();
    switch (tactis) {
        case 0: timer.duration_gpu("Resize(nearest) in gpu takes:"); break;
        case 1: timer.duration_gpu("Resize(bilinear) in gpu takes:"); break;
        case 2: timer.duration_gpu("Resize(bilinear-letterbox) in gpu takes:"); break;
        case 3: timer.duration_gpu("Resize(bilinear-letterbox-center) in gpu takes:"); break;
        default: break;
    }

    // 将结果返回给host上
    CUDA_CHECK(cudaMemcpy(h_tar.data, d_tar, tar_size, cudaMemcpyDeviceToHost));


    CUDA_CHECK(cudaFree(d_src));
    CUDA_CHECK(cudaFree(d_tar));

    return h_tar;
}

这段代码提供了两个函数 preprocess_cpu 和 preprocess_gpu，用于图像预处理，包括缩放和颜色空间转换，分别在 CPU 和 GPU 上执行。下面是对这两个函数的详细分析：(from chatGPT)

CPU 版本预处理 (preprocess_cpu)

• 函数参数：接收一个源图像(cv::Mat &src)、目标高度(const int &tar_h)、目标宽度(const int &tar_w)、一个 Timer 实例和一个操作策略(int tactis)作为参数。返回处理后的图像。
• 缩放逻辑：
  • 首先，计算原图像的高度和宽度比例，然后根据目标高宽比例调整图像大小，确保图像不会因缩放而失真。
  • 使用 cv::resize 函数进行双线性缩放（cv::INTER_LINEAR），将图像缩放到目标大小。
• 颜色空间转换：使用 cv::cvtColor 将图像从 BGR 转换为 RGB 格式。
• 计时功能：通过 Timer 类实例的 start_cpu 和 stop_cpu 方法测量缩放操作的耗时，并通过 duration_cpu 方法打印耗时信息。

GPU 版本预处理 (preprocess_gpu)

• 函数参数：与 CPU 版本类似，但是处理过程在 GPU 上执行。
• 内存分配与拷贝：
  • 使用 cudaMalloc 在 GPU 上为源图像和目标图像分配内存。
  • 通过 cudaMemcpy 将源图像数据从主机复制到 GPU。
• GPU处理：
  • 调用一个自定义的核函数 resize_bilinear_gpu 进行图像的缩放和颜色空间转换，具体实现细节没有在此代码段中展示，但可以推断该函数根据 tactis 参数选择不同的缩放策略。
  • 使用 cudaDeviceSynchronize 确保 GPU 上的所有操作都完成。
• 计时与结果获取：
  • 使用 Timer 类实例的 start_gpu 和 stop_gpu 方法测量 GPU 操作的耗时。
  • 根据 tactis 参数的不同，打印出对应的耗时信息。
  • 最后，使用 cudaMemcpy 将处理后的图像数据从 GPU 复制回主机。
• 资源释放：使用 cudaFree 释放 GPU 上分配的内存。

总结

这两个函数展示了在 CPU 和 GPU 上进行图像预处理的不同方法，包括颜色空间转换和图像缩放。GPU 版本需要显式管理内存（包括分配和释放），并且利用 CUDA 提供的 API 执行数据传输和同步。这种处理方式能够利用 GPU 的并行处理能力，加速图像处理任务，特别是在处理大量数据或进行复杂运算时。

最后我们看下核心代码 preprocess.cu，如下所示：

#include "cuda_runtime_api.h"
#include "stdio.h"
#include <iostream>

#include "utils.hpp"


__global__ void resize_nearest_BGR2RGB_kernel(
    uint8_t* 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 = round((float)y * scaled_h);
    int src_x = round((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) * 3;

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

        // nearest neighbour -- 实现nearest beighbour的resize + BGR2RGB
        tar[tarIdx + 0] = src[srcIdx + 2];
        tar[tarIdx + 1] = src[srcIdx + 1];
        tar[tarIdx + 2] = src[srcIdx + 0];
    }
}

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

    // 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_y1 > srcH || src_x1 > 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) * 3;

        // bilinear interpolation -- 实现bilinear interpolation的resize + 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]);
    }
}

__global__ void resize_bilinear_BGR2RGB_shift_kernel(
    uint8_t* 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_y1 > srcH || src_x1 > 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 -- 计算原图在目标图中的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;

        // 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]);
    }
}

/*
    这里面的所有函数都实现了kernel fusion。这样可以减少kernel launch所产生的overhead
    如果使用了shared memory的话，就可以减少分配shared memory所产生的overhead以及内部线程同步的overhead。(这个案例没有使用shared memory)
    CUDA编程中有一些cuda runtime api是implicit synchronize(隐式同步)的，比如cudaMalloc, cudaMallocHost，以及shared memory的分配。
    高效的CUDA编程需要意识这些implicit synchronize以及其他会产生overhead的地方。比如使用内存复用的方法，让cuda分配完一次memory就一直使用它

    这里建议大家把我写的每一个kernel都拆开成不同的kernel来分别计算
    e.g. resize kernel + BGR2RGB kernel + shift kernel 
    之后用nsight去比较融合与不融合的差别在哪里。去体会一下fusion的好处
*/

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;
    }
    
    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;
    default:
        break;
    }
}

这段代码定义了 GPU 上执行的核心图像处理函数，特别是用于图像缩放和颜色空间转换的 CUDA 核函数。这些函数设计来在 CUDA 架构上高效地处理图像数据。代码中包含了三种核心操作：最近邻插值、双线性插值、以及双线性插值加平移居中。下面详细分析这些组成部分：(from chatGPT)

CUDA 核函数解析

1. 最近邻插值+颜色转换（resize_nearest_BGR2RGB_kernel）

• 功能：对图像进行最近邻插值缩放，并将颜色空间从 BGR 转换为 RGB。
• 参数：包括目标图像和源图像的指针、目标和源图像的宽高、以及缩放因子。
• 实现细节：
  • 计算每个线程应该处理的目标图像上的像素位置（x, y）。
  • 使用缩放因子将目标像素位置映射回源图像上的最近像素。
  • 如果计算出的源像素位置在源图像范围内，进行颜色值的复制和转换，否则不处理越界的部分。
  • 实现了 kernel fusion，即在同一个 CUDA 核函数中完成了缩放和颜色空间转换，减少了内存访问次数。

2. 双线性插值+颜色转换（resize_bilinear_BGR2RGB_kernel）

• 功能：使用双线性插值对图像进行缩放，同时将颜色空间从BGR转换为RGB。
• 参数：同上。
• 实现细节：
  • 计算目标图像上每个像素点对应源图像上的四个近邻像素点位置。
  • 根据四个近邻像素点的颜色值和相对位置，使用双线性插值公式计算目标像素点的颜色值。
  • 在复制和计算颜色值时，同时完成颜色空间的转换。
  • 这种方法可以得到比最近邻插值更平滑的缩放效果。

3. 双线性插值+颜色转换+平移居中（resize_bilinear_BGR2RGB_shift_kernel）

• 功能：在双线性插值和颜色转换的基础上，添加了平移操作以使缩放后的图像在目标空间中居中。
• 参数：同上。
• 实现细节：
  • 首先进行双线性插值和颜色转换。
  • 在将计算得到的颜色值写入目标图像之前，计算目标像素点经过居中处理后的新位置。
  • 如果新位置在目标图像范围内，则将颜色值写入新位置；否则，该像素点不进行处理。

综合函数 resize_bilinear_gpu

• 功能：根据 tactis 参数的值选择执行上述中的一个CUDA核函数。
• 实现细节：
  • 使用 dim3 类型定义了CUDA核函数的网格和块尺寸，确保足够的线程覆盖整个图像。
  • 计算缩放因子，如果 tactis 大于 1，则使用相同的缩放因子，确保图像在缩放时保持比例。
  • 根据 tactis 值选择相应的 CUDA 核函数执行，其中包括最近邻插值、双线性插值，以及双线性插值加平移居中。

总结

这段代码高效地利用 CUDA 进行图像处理，通过并行计算加速了图像缩放和颜色空间转换的操作。代码中的 kernel fusion 技术减少了内存访问次数，提高了处理效率。

OK，以上就是对 2.10-bilinear-interpolation 案例代码的简单分析，最后我们执行下看下输出结果，如下图所示：

其实图片分辨率越大加速比其实越明显，大家可以尝试下不同分辨率的图片看看

总结

本次课程我们主要讲解了双线性插值算法以及如何通过 CUDA 高并发的特性来对其加速，这个核函数的实现大家需要掌握，因为这是我们在 TensorRT 模型部署中非常常见的。博主之前有相关笔记记录过双线性插值因此这里就简单过一下，大家感兴趣的可以多看看案例代码分析分析
OK，以上就是第 5 小节有关双线性插值的全部内容了，下节我们进入第三章—TensorR基础入门的学习，敬请期待😄

参考

• YOLOv5推理详解及预处理高性能实现