TensorRT笔记五（Warpaffine）

• 使用tensorrt进行模型推理前，使用cuda进行图像预处理，可以直接将结果输入tensorrt提高速度，平时目标检测推理的图像预处理主要为warpaffine的缩放和平移，将图像等比缩放并居中。

一、缩放、平移矩阵

1. 缩放

$\left[\begin{array}{c}x\prime \\ y\prime \end{array}\right]$=$\left[\begin{array}{cc}scale& 0\\ 0& scale\end{array}\right]$·$\left[\begin{array}{c}x\\ y\end{array}\right]$

2. 平移

$\left[\begin{array}{c}x\prime \\ y\prime \end{array}\right]$=$\left[\begin{array}{cc}1& 0\\ 0& 1\end{array}\right]$·$\left[\begin{array}{c}x\\ y\end{array}\right]$ + $\left[\begin{array}{c}ox\\ oy\end{array}\right]$

• 合并为一个矩阵变换

$\left[\begin{array}{c}x\prime \\ y\prime \\ w\end{array}\right]$=$\left[\begin{array}{ccc}1& 0& ox\\ 0& 1& oy\\ 0& 0& 1\end{array}\right]$·$\left[\begin{array}{c}x\\ y\\ 1\end{array}\right]$

3.图像等比缩放并居中

分为三个步骤

• 等比缩放
  $$scale=min(\frac{dst.width}{src.width},\frac{dst.height}{src.height})$$
• 平移到src原点(opencv中的图片左上角)
  $$ox=-\frac{scale\times src.width}{2},oy=-\frac{scale\times src.height}{2}$$
• 平移到dst中心位置
  $$ox=\frac{dst.width}{2},oy=\frac{dst.height}{2}$$
• 等比缩放并居中矩阵变换
  $\left[\begin{array}{c}x\prime \\ y\prime \end{array}\right]$=$\left[\begin{array}{ccc}scale& 0& -\frac{scale\times src.width}{2}+\frac{dst.width}{2}\\ 0& scale& -\frac{scale\times src.height}{2}+\frac{dst.height}{2}\end{array}\right]$·$\left[\begin{array}{c}x\\ y\\ 1\end{array}\right]$

3.逆变换

• 在warpaffine实现过程需要其逆变换
  $k=scale$
  $b1=-\frac{scale\times src.width}{2}+\frac{dst.width}{2}$
  $b2=-\frac{scale\times src.height}{2}+\frac{dst.height}{2}$
  $x\prime =kx+b1$
  $y\prime =ky+b2$
  $x=\frac{x\prime -b1}{k}=x\prime \times \frac{1}{k}+(-\frac{b1}{k})$
  $y=\frac{x\prime -b2}{k}=y\prime \times \frac{1}{k}+(-\frac{b2}{k})$
  $W^{-1}$=$\left[\begin{array}{ccc}\frac{1}{k}& 0& -\frac{b1}{k}\\ 0& \frac{1}{k}& -\frac{b2}{k}\end{array}\right]$

二、cuda实现

1. main.cpp

• 使用cudaMalloc在GPU上开辟两块空间&psrc_device和&pdst_device，大小分别为src_size和dst_size
• 将数据搬运到GPU

cudaMemcpy(psrc_device, image.data, src_size, cudaMemcpyHostToDevice);
//image为Mat格式，image.data返回image的首地址

2. warpaffine.cu

• host计算变换逆矩阵
• 开启dst图像大小个线程

 dim3 block_size(32, 32); //blocksize最大就是1024
 dim3 grid_size((dst_width + 31) / 32, (dst_height + 31) / 32);//向上取整

• 核函数warp_affine_bilinear_kernel<<<grid_size, block_size, 0, nullptr>>>
• 线程索引并return多余线程

int dx = blockDim.x * blockIdx.x + threadIdx.x; 
int dy = blockDim.y * blockIdx.y + threadIdx.y;
if (dx >= dst_width || dy >= dst_height)  return;
//(dx,dy)即为dst图像对应的像素

• 由逆矩阵计算dst对应在src上的坐标
• 变换后超出src范围(src_x < -1 || src_x >= src_width || src_y < -1 || src_y >= src_height)的坐标取固定值fill_value
• 其余坐标的值由双线性插值计算
  比如，计算左上角src的值时

uint8_t const_values[] = {fill_value, fill_value, fill_value};
uint8_t* v1 = const_values;
v1 = src + y_low * src_line_size + x_low * 3;
//src 表示存储图像像素值的连续缓冲区的起始地址
//y_low 表示像素所处的行数
//src_line_size 表示缓冲区一行像素值的字节数
//x_low 表示像素所处的列数，
//v1 则表示对应像素的存储地址
//因为一个像素的颜色值通常是由 3 个字节组成，因此需要将 x 坐标乘以 3，以计算出像素在缓冲区中的字节偏移量。

• RGB三个值通过地址的偏移的值的双线性插值计算

c0 = floorf(w1 * v1[0] + w2 * v2[0] + w3 * v3[0] + w4 * v4[0] + 0.5f);
c1 = floorf(w1 * v1[1] + w2 * v2[1] + w3 * v3[1] + w4 * v4[1] + 0.5f);
c2 = floorf(w1 * v1[2] + w2 * v2[2] + w3 * v3[2] + w4 * v4[2] + 0.5f);

3. main.cpp

• 处理完的数据从gpu搬回来

checkRuntime(cudaMemcpy(output.data, pdst_device, dst_size, cudaMemcpyDeviceToHost));

• 释放内存