网站建设宁夏凤凰云,深圳多语言网站建设,做网站在哪里添加关键词,微信商店小程序制作教程这里学习一下定点数的优化操作#xff0c;实际上就是以整数代替浮点数#xff0c;乘除法的操作均通过左右移位来实现#xff0c;适合在算力非常低的场景下使用#xff0c;极致的压榨性能。 https://zhuanlan.zhihu.com/p/338588296 定点数介绍
以下给出函数的具体实现实际上就是以整数代替浮点数乘除法的操作均通过左右移位来实现适合在算力非常低的场景下使用极致的压榨性能。 https://zhuanlan.zhihu.com/p/338588296 定点数介绍
以下给出函数的具体实现函数 convertNV12toYUV444withActions_cuda的作用是 1 把NV12(YUV420uv交替出现)格式的图像转换成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_18numBlocks, 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_cudanumB, threadsPerBlock, 0, stream(nv12buf, srcuvbuf,src_width, src_height);ResizeBilinear_U8_Q13_18numBlocks, 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_18numBlocks, 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_cudaBLOCK_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个u1个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
把uint83通道的数据转成f323通道的数据 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;}
}
