沙坪坝做网站,制作网站一般是多大,中国品牌加盟网官网,个人网站申请备案目录 前言0. 简述1. 执行一下我们的第十个CUDA程序2. Bilinear interpolation3. 代码分析总结参考 前言 自动驾驶之心推出的 《CUDA与TensorRT部署实战课程》#xff0c;链接。记录下个人学习笔记#xff0c;仅供自己参考 Note#xff1a;关于 CUDA 加速双线程插值的内容博主… 目录 前言0. 简述1. 执行一下我们的第十个CUDA程序2. Bilinear interpolation3. 代码分析总结参考 前言 自动驾驶之心推出的 《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 BGR2RGBtar[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) 像素值该怎么求呢我们可以通过加权平均的方式来求取即 d s t [ 100 , 100 ] s r c [ 100 , 200 ] × a 11 s r c [ 101 , 200 ] × a 10 s r c [ 100 , 201 ] × a 01 s r c [ 101 , 201 ] × a 00 dst[100,100] src[100,200]\times a11 src[101,200]\times a10 src[100,201] \times a01 src[101,201] \times a00 dst[100,100]src[100,200]×a11src[101,200]×a10src[100,201]×a01src[101,201]×a00
这个其实就是双线性插值的过程通过未知像素点周围的四个像素来进行加权求和得到最终的像素那依此类推最近邻插值是怎么做的呢那通过名字我们就可以知道这个插值方式其实就是选取距离未知像素点最近的那个像素点的像素值最近邻插值计算速度相较于双线性插值要快但图像质量相较于双线性插值要差
如果我们想要缩放后的图像仍然保持相同的长宽比我们需要保证图像的宽和高的缩放比一致最终的缩放因子 s c a l e m i n ( t a r _ w s r c _ w , t a r _ h s r c _ h ) scale min(\frac{tar\_w}{src\_w}, \frac{tar\_h}{src\_h}) scalemin(src_wtar_w,src_htar_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.hppusing 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_cpuTimer::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 BGR2RGBtar[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 BGR2RGBtar[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 BGR2RGBtar[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 resizefloat 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推理详解及预处理高性能实现
