cuda中的定点数优化技术

本文主要是介绍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 uResizeBilinear_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 vu8c3_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.5val1 = (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;}
}

这篇关于cuda中的定点数优化技术的文章就介绍到这儿，希望我们推荐的文章对编程师们有所帮助！

---