Cuda编程加速图像预处理

文章目录

项目简介
加速效果
Cuda代码
Git地址

项目简介

基于 cuda 和 opencv 环境
目标：
- 单独使用，以加速图像处理操作；
- 结合 TensorRT 使用，进一步加快推理速度

加速效果

这里对比 Deeplabv3+ 使用 cuda 预处理前后的 tensorrt 推理速度
未使用cuda图像预处理的代码，可参考博主的另一个 tensorrt 的项目

FP32精度:

C++图像预处理	cuda图像预处理
25 ms	19 ms

Int8量化后:

C++图像预处理	cuda图像预处理
10 ms	3 ms

Cuda代码

核心的核函数代码如下所示：

BGR to RGB

__global__ void RGB(const uchar* srcData, uchar* tgtData, const int h, const int w)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * w;int idx3 = idx * 3;if (ix < w && iy < h){tgtData[idx3] = srcData[idx3 + 2];tgtData[idx3 + 1] = srcData[idx3 + 1];tgtData[idx3 + 2] = srcData[idx3];}
}

Bilinear resize

__global__ void linear(const uchar* srcData, const int srcH, const int srcW, uchar* tgtData, const int tgtH, const int tgtW)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * tgtW;int idx3 = idx * 3;float scaleY = (float)tgtH / (float)srcH;float scaleX = (float)tgtW / (float)srcW;// (ix,iy)为目标图像坐标// (before_x,before_y)原图坐标float beforeX = float(ix + 0.5) / scaleX - 0.5;float beforeY = float(iy + 0.5) / scaleY - 0.5;// 原图像坐标四个相邻点// 获得变换前最近的四个顶点,取整int topY = static_cast<int>(beforeY);int bottomY = topY + 1;int leftX = static_cast<int>(beforeX);int rightX = leftX + 1;//计算变换前坐标的小数部分float u = beforeX - leftX;float v = beforeY - topY;if (ix < tgtW && iy < tgtH){// 如果计算的原始图像的像素大于真实原始图像尺寸if (topY >= srcH - 1 && leftX >= srcW - 1)  //右下角{for (int k = 0; k < 3; k++){tgtData[idx3 + k] = (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k];}}else if (topY >= srcH - 1)  // 最后一行{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (u) * (1. - v) * srcData[(rightX + topY * srcW) * 3 + k];}}else if (leftX >= srcW - 1)  // 最后一列{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (1. - u) * (v) * srcData[(leftX + bottomY * srcW) * 3 + k];}}else  // 非最后一行或最后一列情况{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (u) * (1. - v) * srcData[(rightX + topY * srcW) * 3 + k]+ (1. - u) * (v) * srcData[(leftX + bottomY * srcW) * 3 + k]+ u * v * srcData[(rightX + bottomY * srcW) * 3 + k];}}}
}

HWC to CHW

__global__ void toCHW(const uchar* srcData, uchar* tgtData, const int h, const int w)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * w;int idx3 = idx * 3;if (ix < w && iy < h){tgtData[idx] = srcData[idx3];tgtData[idx + h * w] = srcData[idx3 + 1];tgtData[idx + h * w * 2] = srcData[idx3 + 2];}
}

Normalize

__global__ void norm(const uchar* srcData, float* tgtData, const int h, const int w)
{/*mean = [0.485, 0.456, 0.406]std = [0.229, 0.224, 0.225](img / 255. - mean) / std*/int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * w;int idx3 = idx * 3;if (ix < w && iy < h){tgtData[idx3] = ((float)srcData[idx3] / 255.0 - 0.406) / 0.225;  // B pixeltgtData[idx3 + 1] = ((float)srcData[idx3 + 1] / 255.0 - 0.456) / 0.224;  // G pixeltgtData[idx3 + 2] = ((float)srcData[idx3 + 2] / 255.0 - 0.485) / 0.229;  // R pixel}
}

综合以上预处理操作（注意：并不是简单的拼接）

// cuda image preprocess__global__ void resize(const uchar* srcData, const int srcH, const int srcW, uchar* tgtData, const int tgtH, const int tgtW)
{int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * tgtW;int idx3 = idx * 3;float scaleY = (float)tgtH / (float)srcH;float scaleX = (float)tgtW / (float)srcW;// (ix,iy)为目标图像坐标// (before_x,before_y)原图坐标float beforeX = float(ix + 0.5) / scaleX - 0.5;float beforeY = float(iy + 0.5) / scaleY - 0.5;// 原图像坐标四个相邻点// 获得变换前最近的四个顶点,取整int topY = static_cast<int>(beforeY);int bottomY = topY + 1;int leftX = static_cast<int>(beforeX);int rightX = leftX + 1;//计算变换前坐标的小数部分float u = beforeX - leftX;float v = beforeY - topY;if (ix < tgtW && iy < tgtH){// 如果计算的原始图像的像素大于真实原始图像尺寸if (topY >= srcH - 1 && leftX >= srcW - 1)  //右下角{for (int k = 0; k < 3; k++){tgtData[idx3 + k] = (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k];}}else if (topY >= srcH - 1)  // 最后一行{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (u) * (1. - v) * srcData[(rightX + topY * srcW) * 3 + k];}}else if (leftX >= srcW - 1)  // 最后一列{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (1. - u) * (v) * srcData[(leftX + bottomY * srcW) * 3 + k];}}else  // 非最后一行或最后一列情况{for (int k = 0; k < 3; k++){tgtData[idx3 + k]= (1. - u) * (1. - v) * srcData[(leftX + topY * srcW) * 3 + k]+ (u) * (1. - v) * srcData[(rightX + topY * srcW) * 3 + k]+ (1. - u) * (v) * srcData[(leftX + bottomY * srcW) * 3 + k]+ u * v * srcData[(rightX + bottomY * srcW) * 3 + k];}}}
}__global__ void process(const uchar* srcData, float* tgtData, const int h, const int w)
{/*mean = [0.485, 0.456, 0.406]std = [0.229, 0.224, 0.225](img / 255. - mean) / std*/int ix = threadIdx.x + blockIdx.x * blockDim.x;int iy = threadIdx.y + blockIdx.y * blockDim.y;int idx = ix + iy * w;int idx3 = idx * 3;if (ix < w && iy < h){tgtData[idx] = ((float)srcData[idx3 + 2] / 255.0 - 0.485) / 0.229;  // R pixeltgtData[idx + h * w] = ((float)srcData[idx3 + 1] / 255.0 - 0.456) / 0.224;  // G pixeltgtData[idx + h * w * 2] = ((float)srcData[idx3] / 255.0 - 0.406) / 0.225;  // B pixel}
}void preprocess(const cv::Mat& srcImg, float* dstData, const int dstHeight, const int dstWidth)
{int srcHeight = srcImg.rows;int srcWidth = srcImg.cols;int srcElements = srcHeight * srcWidth * 3;int dstElements = dstHeight * dstWidth * 3;// target data on devicefloat* dstDevData;cudaMalloc((void**)&dstDevData, sizeof(float) * dstElements);// middle image data on device ( for bilinear resize )uchar* midDevData;cudaMalloc((void**)&midDevData, sizeof(uchar) * dstElements);// source images data on deviceuchar* srcDevData;cudaMalloc((void**)&srcDevData, sizeof(uchar) * srcElements);cudaMemcpy(srcDevData, srcImg.data, sizeof(uchar) * srcElements, cudaMemcpyHostToDevice);dim3 blockSize(32, 32);dim3 gridSize((dstWidth + blockSize.x - 1) / blockSize.x, (dstHeight + blockSize.y - 1) / blockSize.y);// bilinear resizeresize<<<gridSize, blockSize>>>(srcDevData, srcHeight, srcWidth, midDevData, dstHeight, dstWidth);cudaDeviceSynchronize();// hwc to chw / bgr to rgb / normalizeprocess<<<gridSize, blockSize>>>(midDevData, dstDevData, dstHeight, dstWidth);cudaMemcpy(dstData, dstDevData, sizeof(float) * dstElements, cudaMemcpyDeviceToHost);cudaFree(srcDevData);cudaFree(midDevData);cudaFree(dstDevData);
}

Git地址

完整的cuda图像预处理代码链接：cuda-image-preprocess

Cuda编程加速图像预处理相关推荐

cuda编程思想和opencv_gpu图像处理
CUDA编程 https://github.com/Ewenwan/ShiYanLou/tree/master/CUDA CUDA(Compute Unified Device Architectur ...
CUDA编程之快速入门-----GPU加速原理和编程实现
转载:https://www.cnblogs.com/skyfsm/p/9673960.html CUDA(Compute Unified Device Architecture)的中文全称为计算统一 ...
【CV学习笔记】图像预处理warpaffine-cuda加速
1.前言在上个学习笔记中学习warpaffine,并且在opencv下面实现了图像的预处理,而warpaffine可以很好的利用cuda加速来实现,于是基于手写AI的项目,又学习了warpaffie ...
Yolov3(Mxnet)测试加速：GPU图像预处理
深度学习的数据标准化操作在测试的也要遵守,但是Mxnet中Gluoncv使用CPU的串行数据标准化,对于某些实时性要求较高的任务,在CPU使用率较高时,数据标准化的耗时严重拖累了网络的预测速度.我们以 ...
GPU — CUDA 编程模型
目录文章目录目录 GPGPU CUDA 编程模型 CUDA 的架构 CUDA 的工作原理 Grid.Block.Thread Warp GPGPU GPU 起初是用来处理图像的,但是后来人们发现其 ...
linux运行并行计算cuda,并行化计算与CUDA编程
原标题:并行化计算与CUDA编程近年来,显卡的更新换代也比较快,NVIDIA今年的发布会也宣布了RTX3080TI即将到来.显卡的运算能力也越来越强. 很多人对显卡的了解可能源于游戏,但是你的显卡不 ...
CUDA编程：与OpenCV结合
前言学习计算机图像处理算法的童鞋,就不得不学习cuda,为啥呢?因为图像处理一般都是矩阵运算,动不动就是百万的计算量这个时候优化计算时间是必不可少的.openCV本身提供了很多cuda函数,能够满足 ...
CUDA 编程学习
0 简单的CUDA简介 1.简单教程 CUDA C ++只是使用CUDA创建大规模并行应用程序的方法之一.它允许您使用功能强大的C ++编程语言来开发由GPU上运行的数千个并行线程加速的高性能算法.许 ...
cuda学习(5):使用cuda核函数加速warpaffine
1. warpaffine 介绍放射变换(warpaffine),主要解决图像的缩放和平移来处理目标检测中常见的预处理行为. 比如有一张猫的图片,对图片进行letterbox变换,将图片进行缩放,并 ...

Cuda编程加速图像预处理

文章目录

项目简介

加速效果

Cuda代码

Git地址

Cuda编程加速图像预处理相关推荐

最新文章

热门文章