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;}
}

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/news/639181.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

纸黄金实战投资技巧:避免亏损的有效策略

在纸黄金交易的实战中&#xff0c;避免亏损是每位投资者都追求的目标。虽然任何投资都存在一定的风险&#xff0c;但采取一些有效的策略可以帮助投资者最大限度地减少亏损的可能性。以下是一些在纸黄金交易中避免亏损的实战技巧&#xff1a; 一、设定止损点是避免亏损的关键 止…

【Linux】

Linux零基础入门 列出文件/文件夹新建/切换路径查看当前路径重命名或者移动文件夹拷贝文件/文件夹删除文件夹设置环境变量编辑文本文件压缩和解压查看cpu的信息查看/杀死进程查看进程的CPU和内存占用重定向日志场景一场景二场景三场景四 列出文件/文件夹 命令&#xff1a;Ls(L…

All the stories begin at installation

Before installation, there are some key points about Conan: “Conan is a dependency and package manager for C and C languages.”“With full binary management, Conan can create and reuse any number of different binaries (for different configurations like a…

基于SpringBoot的智慧社区居家养老健康管理系统

文章目录 项目介绍主要功能截图&#xff1a;部分代码展示设计总结项目获取方式 &#x1f345; 作者主页&#xff1a;超级无敌暴龙战士塔塔开 &#x1f345; 简介&#xff1a;Java领域优质创作者&#x1f3c6;、 简历模板、学习资料、面试题库【关注我&#xff0c;都给你】 &…

Kafka常见指令及监控程序介绍

kafka在流数据、IO削峰上非常有用&#xff0c;以下对于这款程序&#xff0c;做一些常见指令介绍。 下文使用–bootstrap-server 10.0.0.102:9092,10.0.0.103:9092,10.0.0.104:9092 需自行填写各自对应的集群IP和kafka的端口。 该写法 等同 –bootstrap-server localhost:9092 …

Unity 抽象工厂模式(实例详解)

文章目录 简介实例1实例2 简介 抽象工厂模式是一种创建型设计模式&#xff0c;它提供了一种方式来封装一组相关或相互依赖对象的创建过程&#xff0c;而无需指定具体类。这种模式常用于系统中有多组相关产品族&#xff0c;且客户端需要使用不同产品族中的对象时。 在Unity中&a…

Windows7关闭谷歌浏览器提示“若要接收后续 Google Chrome 更新,您需使用 Windows 10 或更高版本”的方法

背景 电脑比较老&#xff0c;系统一直没有更新&#xff0c;硬件和软件版本如下&#xff1a; 操作系统版本&#xff1a;Windows7 企业版 谷歌浏览器版本&#xff1a;109.0.5414.120&#xff08;正式版本&#xff09; &#xff08;64 位&#xff09; 该版本的谷歌浏览器是支持…

MySQL不同插入方式性能对比实验

最近负责的项目需要数据同步入库MySQL&#xff0c;为了测速那种入库方式效率比较高&#xff0c;为此进行了以下的对比实验&#xff0c;在此记录一下 实验表单数据格式 实验代码 共三种方法对比 mutiSqlInsert: 一条一条插入&#xff0c;最后一次提交 singleSqlInsert&…

LabVIEW滚动轴承故障在线监测

展示了如何将LabVIEW开发出一种有效的滚动轴承故障在线监测系统。介绍了该系统的开发过程、工作原理及其在实际应用中的效果。该系统成功地应用于对滚动轴承故障的早期诊断&#xff0c;提高了故障检测的准确性和效率。 滚动轴承在工作过程中会产生复杂的振动信号&#xff0c;包…

Proxmox VE 8 试装Oracle 23c

作者&#xff1a;田逸&#xff08;formyz&#xff09; Oracle 当前的最新版本是23c&#xff0c;虽然官方网站下载不了它的正式版本&#xff0c;但是却提供了一个性能受限的免费版本“Oracle Database 23.3 Free”&#xff08;存储容量受限、内存使用受限&#xff09;。这里就只…

[机缘参悟-129] :个人对人生之苦解决之道的思考

目录 前言&#xff1a; 第1层&#xff1a;环境层 1.1 环境的分类 1.2 理解环境的运作的基本原理 1.3 主动选择适合自己的愉快的环境 1.4 主动构建适合自己的愉快的环境 第2层&#xff1a;生理层 2.1 生理健康和情绪之间的关系 2.2 学习人的生物、生理、健康的基本知识…

MySQL与PostgreSQL对比

对比 许可证 License MySQL 社区版采用 GPL 许可证。Postgres 发布在 PostgreSQL 许可下&#xff0c;是一种类似于 BSD 或 MIT 的自由开源许可。 即便 MySQL 采用了 GPL&#xff0c;仍有人担心 MySQL 归 Oracle 所有&#xff0c;这也是为什么 MariaDB 从 MySQL 分叉出来。 …

数据出境——电商API接口使得电商数据跨境流动已成为趋势

在数字经济的当下&#xff0c;数据已经成为商业决策的关键因素。尤其是电商领域&#xff0c;电商API数据采集不仅关乎企业运营效率&#xff0c;还涉及到用户隐私与国家安全。近年来&#xff0c;随着电商市场的全球化发展&#xff0c;电商数据出境与跨境贸易已成为不可逆转的趋势…

vue3-模版引用ref

1. 介绍 概念&#xff1a;通过 ref标识 获取真实的 dom对象或者组件实例对象 2. 基本使用 实现步骤&#xff1a; 调用ref函数生成一个ref对象 通过ref标识绑定ref对象到标签 代码如下&#xff1a; 父组件&#xff1a; <script setup> import { onMounted, ref } …

Android Studio 之 菜单 Menu

选项菜单 OptionsMenu 用xml添加&#xff08;更建议使用&#xff09; 创建一个菜单布局 : 在 res文件下新建一个menu 目录&#xff0c;此时的菜单id为&#xff1a;R.menu.option <?xml version"1.0" encoding"utf-8"?> <menu xmlns:android&…

不同开发语言在进程、线程和协程的设计差异

不同开发语言在进程、线程和协程的设计差异 1. 进程、线程和协程上的差异1.1 进程、线程、协程的定义1.2 进程、线程、协程的差异1.3 进程、线程、协程的内存成本1.4 进程、线程、协程的切换成本 2. 线程、协程之间的通信和协作方式2.1 python如何实现线程通信&#xff1f;2.2 …

【Unity】AB包下载

【Unity】AB包下载 1.使用插件打AB包 a.AB包分类 一般地&#xff0c;将预制体作为AB包资源&#xff0c;不仅需要对预制体本身进行归类&#xff0c;还要对其涉及的动画&#xff08;AnimationClip&#xff09;、动画状态机&#xff08;AnimatorController&#xff09;、以及所…

《A++ 敏捷开发》- 5 量化管理从个人开始

我&#xff1a;你们管理层和客户都比较关心项目的进度&#xff0c;项目是否能按时完成&#xff1f;请问你们过去的项目如何&#xff1f; 开发&#xff1a;我们现在就是走敏捷开发&#xff0c;两周一个迭代。每次迭代前&#xff0c;我们聚一起开会&#xff0c;把所有用户故事按优…

Dubbo 3.2版本分析Provider启动时操作

Dubbo 3.2版本分析Provider启动时操作 前言例子分析onStarting 模块doStart 模块 小结 前言 上一篇文章&#xff0c;我们分析了 Dubbo 3.2 版本在 Provider 启动前的操作流程&#xff0c;这次我们具体分析具体它的启动过程&#xff0c;揭开它的神秘面纱。 例子 这里我们还是…

【ZYNQ入门】第八篇、基于Lwip构建TCP服务器

目录 第一部分、基础知识 1、小白入门必看文章 2、什么是Lwip&#xff1f; 3、什么是TCP/IP协议&#xff1f; 4、MAC地址、IP地址、子网掩码、网关 4.1、MAC地址 4.2、IP地址 4.3、子网掩码 4.4、网关 第二部分、硬件搭建 第三部分、软件代码 1、SDK工程的建立 2、…