cuda中的定点数优化技术

这里学习一下定点数的优化操作，实际上就是以整数代替浮点数，乘除法的操作均通过左右移位来实现，适合在算力非常低的场景下使用，极致的压榨性能。
https://zhuanlan.zhihu.com/p/338588296 定点数介绍

以下给出函数的具体实现，函数 convertNV12toYUV444withActions_cuda的作用是
1 把NV12(YUV420，uv交替出现)格式的图像转换成YUV444的格式
2 进行crop/resize的算子操作
3 同时对数据做toFloat操作

具体实现过程是，把Y数据经过ResizeBilinear_U8_Q13_18放入tmpImagecuda中，紧跟通过nv12_extract_uv44_cuda提取UV的数据，再次使用ResizeBilinear_U8_Q13_18放入tmpImagecuda后续空间中，最后u8c3_convertTo_f32c3_cuda 完成toFloat的操作，放入dst_img中

void convertNV12toYUV444withActions_cuda(uint8_t *src_img,
										 uint8_t *src_imgcuda,
                                         uint8_t *tmpImagecuda,
                                         ImageTransParam &trans_param,
                                         uint8_t *dst_imgcuda,
                                         uint8_t *dst_img, cudaStream_t stream) {
  int src_width = trans_param.src_width;
  int src_height = trans_param.src_height;
  int resize_width = trans_param.resize_width;
  int resize_height = trans_param.resize_height;
  int crop_width = resize_width;
  int crop_height = resize_height;
  int xscale = (src_width << 18) / resize_width;
  int yscale = (src_height << 18) / resize_height;
  int crop_start_x = 0;
  int crop_start_y = 0;
  if (trans_param.is_crop) {
    crop_width = trans_param.crop_width;
    crop_height = trans_param.crop_height;
    crop_start_x = (trans_param.crop_start_x * xscale) >> 18;
    crop_start_y = (trans_param.crop_start_y * yscale) >> 18;
  }
  int srclen = src_width * src_height;
  int dstlen = crop_width * crop_height;
  uint8_t *nv12buf;
  uint8_t *cropuvbuf;
  uint8_t *srcuvbuf;

  int size1 =
      (srclen * 2) + dstlen * 3;  // tmpImage size
                                  //+dstlen* 3 *sizeof(float) +srclen*1.5;

  int size_dst = dstlen * 3 * sizeof(float);
  src_imgcuda = src_img;
  // cudaMemcpy(src_imgcuda, src_img, srclen*3/2, cudaMemcpyHostToDevice);

  dim3 threadsPerBlock(16, 16);
  dim3 numBlocks((crop_width + threadsPerBlock.x - 1) / (threadsPerBlock.x),
                 (crop_height + threadsPerBlock.y - 1) / (threadsPerBlock.y));
  ResizeBilinear_U8_Q13_18<<<numBlocks, threadsPerBlock, 0 , stream>>>(
      &src_imgcuda[crop_start_y * src_width + crop_start_x], tmpImagecuda,
      src_width, crop_width, crop_width, crop_height, xscale, yscale, 0, 0);

  nv12buf = &src_imgcuda[srclen];
  cropuvbuf = &tmpImagecuda[dstlen];
  srcuvbuf = &tmpImagecuda[dstlen * 3];

  dim3 numB(((src_width >> 1) + threadsPerBlock.x - 1) / (threadsPerBlock.x),
            ((src_height >> 1) + threadsPerBlock.y - 1) / (threadsPerBlock.y));
  nv12_extract_uv44_cuda<<<numB, threadsPerBlock, 0, stream>>>(nv12buf, srcuvbuf,
                                                    src_width, src_height);

  ResizeBilinear_U8_Q13_18<<<numBlocks, threadsPerBlock, 0, stream>>>(
      &srcuvbuf[crop_start_y * src_width + crop_start_x], cropuvbuf, src_width,
      crop_width, crop_width, crop_height, xscale, yscale, 0, 0);  // resize u
  ResizeBilinear_U8_Q13_18<<<numBlocks, threadsPerBlock, 0 ,stream>>>(
      &srcuvbuf[crop_start_y * src_width + crop_start_x + srclen],
      &cropuvbuf[dstlen], src_width, crop_width, crop_width, crop_height,
      xscale, yscale, 0, 0);  // resize v

  u8c3_convertTo_f32c3_cuda<<<BLOCK_NUM, THREAD_NUM, 0, stream>>>(
      tmpImagecuda, (float *)dst_img, 1 / 128.0, -1.0,
      dstlen * 3);  // dst_imgcuda

  // cudaDeviceSynchronize();
  // cudaMemcpy((void*)dst_img, (void*)dst_imgcuda, dstlen*3*sizeof(float),
  // cudaMemcpyDeviceToHost);
}

ResizeBilinear_U8_Q13_18

双线性插值的一部分，用于计算目标像素在源图像中对应位置的 y 坐标。y 是根据缩放比例 yscale 和偏移量 ys 计算出的固定点数值。y0 是计算出的整数部分，而 wy 是小数部分，即权重用于插值计算

__global__ void ResizeBilinear_U8_Q13_18(uint8_t *Src, uint8_t *Dst,
                                         int32_t SrcPitch, int32_t DstPitch,
                                         int32_t dstWidth, int32_t dstHeight,
                                         int xscale, int yscale,
                                         int32_t x_offset, int32_t y_offset) {
  int xs = x_offset;
  int ys = y_offset;
  int xd = 0;
  int yd = 0;

  uint8_t *psrc = (uint8_t *)Src;
  uint8_t *pdst = (uint8_t *)Dst;

  int sstride = SrcPitch;
  int dstride = DstPitch;

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

  // for (i = 0; i < dstHeight; ++i)
  if (i < dstHeight) {
    int y = (yd + i) * yscale - (ys << 18) + (yscale >> 1) - (1 << 17);
    int y0 = y >> 18;
    int wy = y & ((1 << 18) - 1);

    // for (j = 0; j < dstWidth; ++j)
    if (j < dstWidth) {
      int x = (xd + j) * xscale - (xs << 18) + (xscale >> 1) - (1 << 17);
      int x0 = x >> 18;
      int wx = x & ((1 << 18) - 1);

      int val0 = psrc[(y0 + 0) * sstride + x0] * ((1 << 18) - wx) +
                 psrc[(y0 + 0) * sstride + x0 + 1] * wx;
      int val1 = psrc[(y0 + 1) * sstride + x0] * ((1 << 18) - wx) +
                 psrc[(y0 + 1) * sstride + x0 + 1] * wx;

      val0 = (val0 + (1 << 12)) >> 13;  // round to Q8.5
      val1 = (val1 + (1 << 12)) >> 13;

      pdst[i * dstride + j] =
          (val0 * ((1 << 18) - wy) + val1 * wy + (1 << 22)) >> 23;
    }
  }
}

双线性插值 的公式和原理

双线性插值（Bilinear Interpolation）是一种在二维空间进行插值的方法，常用于图像处理中的缩放和旋转等操作。它通过在两个方向（通常是水平和垂直）上进行线性插值来计算新的像素值。双线性插值比简单的最邻近插值（nearest-neighbor interpolation）更平滑，但计算上也更为复杂一些。

原理：
双线性插值的基本原理是对四个最接近的像素点（通常是一个像素的左上、右上、左下、右下邻居）进行加权平均。权重取决于目标像素点与这四个邻居的相对距离。
公式：

在图像处理中，双线性插值用于在改变图像大小时计算新像素位置的颜色值。这种方法在缩放图像时可以获得比最邻近插值更平滑的结果，尤其是在放大图像时更为明显。然而，它可能引入一些模糊，特别是在大幅度缩放时。对于更高质量的图像缩放，可以考虑使用双三次插值（Bicubic Interpolation）等更高级的方法。

(xscale >> 1) - (1 << 17)

这里的操作是为了向下取整，好比如果dst中的i像素对应原图中的15.5 像素这里要取到左边的15和右边的16

val0 = (val0 + (1 << 12)) >> 13;  // round to Q8.5

这个操作是将之前计算的固定点数的结果 val0 转换回普通的整数格式。这个过程包括两个步骤：舍入（rounding）和位移（shifting）。让我解释一下这两个步骤是如何工作的：

舍入：

(1 << 12) 产生了一个数值为 4096 的数，这是用于舍入的值。在固定点数学中，添加一半的量级（在这个场景中是 4096，即 2^12）是一个常用的舍入技术。这相当于加上 0.5（在相应的量级下）然后向下取整，从而实现标准的四舍五入。
位移：

13 是将结果右移 13 位。这个操作实际上是在将数值除以 2^13（或 8192）。由于之前在计算过程中数值被放大了（左移操作），现在需要通过右移来将其缩小回原来的量级。
这个组合操作（先加上一个舍入值，然后右移）实际上是将固定点数转换回普通整数的过程。这个过程保留了原始数值的整数部分，同时考虑了小数部分对整数部分的影响（即舍入效果）。

在图像处理和其他需要高性能计算的领域中，使用固定点数学而不是浮点数学是一种常见的优化手段。这样可以在不牺牲太多精度的情况下，提高计算效率。

nv12_extract_uv44_cuda

这里并不涉及定点数，就是提取yuv，提取的策略是，因为yuv是420要提升到444，所以1个u对应提升后的4个u，1个v对应提升后的四个v

__global__ void nv12_extract_uv44_cuda(uint8_t *uv_data, uint8_t *dst,
                                       int width, int height) {
  int len = width * height;
  uint8_t *dstu = dst;
  uint8_t *dstv = dst + len;

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

  // for (int y = 0; y < height >> 1; y++)
  if (y < (height >> 1)) {
    // for (int x = 0; x < width >> 1; x++)
    if (x < (width >> 1)) {
      int uv_index = y * width + x * 2;

      uint8_t color_u = uv_data[uv_index];
      uint8_t color_v = uv_data[uv_index + 1];

      dstu[(2 * y + 0) * width + (2 * x + 0)] = color_u;
      dstu[(2 * y + 0) * width + (2 * x + 1)] = color_u;
      dstu[(2 * y + 1) * width + (2 * x + 0)] = color_u;
      dstu[(2 * y + 1) * width + (2 * x + 1)] = color_u;

      dstv[(2 * y + 0) * width + (2 * x + 0)] = color_v;
      dstv[(2 * y + 0) * width + (2 * x + 1)] = color_v;
      dstv[(2 * y + 1) * width + (2 * x + 0)] = color_v;
      dstv[(2 * y + 1) * width + (2 * x + 1)] = color_v;
    }
  }
}

u8c3_convertTo_f32c3_cuda

把uint8，3通道的数据转成f32，3通道的数据
i用来确定像素在hw中的位置，blockDim.x * gridDim.x用来确定hw的总长度，即stride步长

__global__ void u8c3_convertTo_f32c3_cuda(uint8_t *src, float *dst, float alpha,
                                          float beta, int len) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  for (; i < len; i += blockDim.x * gridDim.x) {
    dst[i] = alpha * src[i] + beta;
  }
}