视频拼接融合产品的产品与架构设计(三)内存和显存单元数据迁移

上一篇文章
视频拼接融合产品的产品与架构设计(二)

这一篇沉下先来,彻底放弃了界面,界面最终的体现是最后要做的,现在要做的是产品的架构,使用链式架构方式迁移数据。同时增加插件口,方便编程序。

插件架构

为了视频拼接和算法等的产品化,在视频解码前(录像),解码后在gpu,解码后转颜色空间(bgr),解码后算法处理,解码后算法处理下放部分数据到cpu(如截图),解码后算法处理后转颜色空间,框架必须能够在各个点上有接口,下放给使用着。其中就包含了内存和显存单元的数据迁移。

目的

先举个简单的例子说明,说明拼接的过程,实际上,拼接就是个算法处理,属于解码后转颜色空间之后的算法部分,其他如AI,去雾,超分算法等等都将在这个口以后进行。

  使用两个或者多个gpu拼接两幅图像以后,往下传,到内存合并,然后转码成nv12,编码,rtsp向外传流。虽然数据下放到了cpu,但是在没有nvlink的显卡上,我们的数据无法直接迁移,只能放到cpu上,这一部分是系统中的例外。如果直接有nvlink,可以通过pcie的数据通道直接从一个显卡到另外一个显卡,那就没有问题,我们保持最差系统配置,没有link接口,多块显卡,数据通过cpu迁移到内存进行处理。

  下面来看nv12 的数据,如下所示,uv 紧密排列,宽度和y分量保持一致,而高度为y分量的一半。

在这里插入图片描述

做法1

一开始使用bgr 往下传,数据量很大,往下传的时候速度也很慢,编码的时候要使用ffmpeg swscale转码bgr到yuv420,最后转成nv12,进行编码,非常慢,因为图像两边都近4k,这个方案pass掉。

做法2

使用unified 内存,cuda中有一个cudaMallocManaged函数来控制内存端和服务端,自动会迁移数据,不用显式拷贝,我们来尝试以下,下面的代码只是申请内存,发现申请unified内存以后,内存和cuda kernel函数同时可以操作,确实省事了不少,随之而来的是更大的问题,操作速度非常慢

#include <cuda_runtime.h>size_t size = width * height * sizeof(float); // 假设为浮点型数据
float* devPtr;cudaError_t err = cudaMallocManaged(&devPtr, size, cudaMemAttachGlobal);
if (err != cudaSuccess) {// 处理错误
}// 现在devPtr可以被CPU和GPU代码直接访问
// 注意:实际使用中还需考虑同步问题,确保数据访问的安全性

经过测试,速度非常慢,比bgr更慢,内存自动迁移对于gpu和cpu来说不适合,还是要手动来做数据迁移

改进1
使用cuda转码bgr到nv12,把nv12往下传,每隔一行拷贝一行,使用异步

void mergeYUV(NV12Image* s1, NV12Image* s2, AVFrame* d) {
//#pragma omp parallel for num_threads(2)cudaStream_t stream;cudaStreamCreate(&stream);for (int i = 0; i < s1->height ; ++i) { // UV高度是Y高度的一半cudaMemcpyAsync(d->data[0] + i * d->linesize[0], s1->Y + i * s1->width, s1->width, cudaMemcpyDeviceToHost, stream);cudaMemcpyAsync(d->data[0] + i * d->linesize[0] + s1->width, s2->Y + i * s2->width, s2->width, cudaMemcpyDeviceToHost, stream);if (i < s1->height / 2){cudaMemcpyAsync(d->data[1] + i * d->linesize[1], s1->UV + i * (s1->width), s1->width, cudaMemcpyDeviceToHost, stream);cudaMemcpyAsync(d->data[1] + i * d->linesize[1] + s1->width, s2->UV + i * (s2->width), s2->width, cudaMemcpyDeviceToHost, stream);}//memcpy(d->data[1] + i * d->linesize[1], s1->UV + i * (s1->width) , s1->width);//memcpy(d->data[1] + i * d->linesize[1] + s1->width , s2->UV + i * (s2->width) , s2->width);}cudaStreamSynchronize(stream);cudaStreamDestroy(stream);
}

这个方法很慢,是因为图像很大,但是比起前面的速度快,整个一帧包含下放和拷贝需要58毫秒,查寻cuda文档后,说明下放次数一定要少,确定以后有信心了,重新组织数据,cuda转空间颜色的函数把nv12放在一块显存中,然后在cudaMemcpy下放迁移。

改进2
以下为实现,算法框架搭建好,开始整理算法,先把两块数据通过迁移直接下放到内存,

int MergeNV12FromGPUMat(cv::cuda::GpuMat& g1, cv::cuda::GpuMat& g2,AVFrame* frame)
{int w1 = g1.cols;int h1 = g1.rows;int w2 = g2.cols;int h2 = g2.rows;int d1YSize = w1 * h1;int d1UVSize = w1 * h1 / 2;int d2YSize = w2 * h2;int d2UVSize = w2 * h2 / 2;void* GPUPtr1 = NULL;void* GPUPtr2 = NULL;double startTime = cv::getTickCount();cudaError_t err1 = cudaMalloc(&GPUPtr1, d1YSize + d1UVSize);if (err1 != cudaSuccess) {return -1;}cudaError_t err2 = cudaMalloc(&GPUPtr2, d2YSize + d2UVSize);if (err2 != cudaSuccess) {cudaFree(&GPUPtr1);return -1;}unsigned char* Y1 = (unsigned char*)GPUPtr1;unsigned char* UV1 = (unsigned char*)GPUPtr1 + w1 * h1;//显卡1bgr24_to_nv12_cuda(reinterpret_cast<uchar3*>(g1.data), Y1,UV1, w1, h1, g1.step, w1, w1);unsigned char* Y2 = (unsigned char*)GPUPtr2;unsigned char* UV2 = (unsigned char*)GPUPtr2 + w2 * h2;
//显卡2    bgr24_to_nv12_cuda(reinterpret_cast<uchar3*>(g2.data), Y2,UV2, w2, h2, g2.step, w2, w2);NV12Image gpum1(Y1, UV1, w1, h1);NV12Image gpum2(Y2, UV2, w2, h2);double endTime = cv::getTickCount();// 计算时间差,单位为毫秒double elapsedTimeMs = (endTime - startTime) / cv::getTickFrequency() * 1000.0;std::cout << "Elapsed time cuda trans in milliseconds: " << elapsedTimeMs << std::endl;// mergeYUV(&gpum1, &gpum2, frame);mergeYUVHost(&gpum1, &gpum2, frame);double endTime1 = cv::getTickCount();elapsedTimeMs = (endTime1 - endTime) / cv::getTickFrequency() * 1000.0;std::cout << "Elapsed time Merge in milliseconds: " << elapsedTimeMs << std::endl;cudaFree(GPUPtr1);cudaFree(GPUPtr2);
}

以下转拷贝计算

    //以下使用cpu做
{void* host1YUV = NULL;void* host2YUV = NULL;int n1 = in1->width * in1->height * 1.5;int n2 = in2->width * in2->height * 1.5;cudaMallocHost(&host1YUV, n1);cudaMallocHost(&host2YUV, n2);cudaMemcpy(host1YUV,in1->Y, n1,cudaMemcpyDeviceToHost);cudaMemcpy(host2YUV,in2->Y, n2,cudaMemcpyDeviceToHost);uint8_t* y1 = (uint8_t*)host1YUV;uint8_t* uv1 = (uint8_t*)host1YUV + in1->width * in1->height;uint8_t* y2 = (uint8_t*)host2YUV;uint8_t* uv2 = (uint8_t*)host2YUV + in2->width * in2->height;NV12Image ss1(y1, uv1, in1->width, in1->height);NV12Image ss2(y2, uv2, in2->width, in2->height);NV12Image* s1 = &ss1;NV12Image* s2 = &ss2;for (int i = 0; i < s1->height; ++i) { // UV高度是Y高度的一半memcpy(d->data[0] + i * d->linesize[0], s1->Y + i * s1->width, s1->width);memcpy(d->data[0] + i * d->linesize[0] + s1->width, s2->Y + i * s2->width, s2->width);if (i < s1->height / 2){memcpy(d->data[1] + i * d->linesize[1], s1->UV + i * (s1->width), s1->width);memcpy(d->data[1] + i * d->linesize[1] + s1->width, s2->UV + i * (s2->width), s2->width);}}//end:cudaFree(host1YUV);cudaFree(host2YUV);
}

看时间,两次转颜色空间由于使用了cuda,速度还行,花费3毫秒不到,我的显卡特意选了16系列,1650的显卡,速度不是很快,但也不至于拖后腿。总共下来花费18毫秒多,比起50毫秒有3倍左右提升,这个还是有点多,仔细思考了一下,想从拷贝上下功夫。
在这里插入图片描述

改进3

改成两个线程启动,一个线程拷贝Y分量,一个线程拷贝uv分量

// 合并Y分量的线程函数
void mergeY(NV12Image* s1,NV12Image* s2, AVFrame* d) 
{
#pragma omp parallel for //num_threads(10)for (int i = 0; i < s1->height; ++i) {memcpy(d->data[0] + i * d->linesize[0], s1->Y + i * s1->width, s1->width);memcpy(d->data[0] + i * d->linesize[0] + s1->width, s2->Y + i * s2->width, s2->width);}
}// 合并UV分量的线程函数
void mergeUV(NV12Image* s1, NV12Image* s2, AVFrame* d) {
#pragma omp parallel for //num_threads(10)for (int i = 0; i < s1->height / 2; ++i) { // UV高度是Y高度的一半memcpy(d->data[1] + i * d->linesize[1], s1->UV + i * (s1->width) , s1->width);memcpy(d->data[1] + i * d->linesize[1] + s1->width , s2->UV + i * (s2->width) , s2->width);}
}

在这里插入图片描述

#if 1std::thread yThread(mergeY,s1,s2, d);std::thread uvThread(mergeUV, s1, s2, d); 等待线程完成yThread.join();uvThread.join();#endif

启动线程做,同时启动openmp,大约有2毫秒的提升,线程启动和释放还是需要时间,这个优化必须让线程一直启动,也不要停止

改进4

启用多线程,同时启用更好的avx2 指令集,具体看我的其他文章,最终锁定在10毫秒以内,包括了转颜色空间,数据迁移到内存,合并成ffmpeg可以编码的nv12格式,花费10毫秒,整个过程比完全在显存里面多了一次数据迁移的过程,就是数据既下载又上传,因为在编码的时候,数据又上传到了gpu,最后推流时,又迁移到了内存,rtsp发送数据流,是一定要到内存里面的。
在这里插入图片描述

总结

下一篇会往整体架构上写,多画图,少代码

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

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

相关文章

短剧奔向小程序,流量生意如何开启?

随着移动互联网的飞速发展&#xff0c;小程序作为一种轻量级、易传播的应用形态&#xff0c;逐渐在各个领域展现出其独特的商业价值。而最近爆火的短剧小视频作为一种受众广泛的娱乐形式&#xff0c;与小程序结合后&#xff0c;不仅为观众提供了更为便捷的观看体验&#xff0c;…

Linux线程(三)死锁与线程同步

目录 一、什么是死锁 死锁的四个必要条件 如何避免死锁 避免死锁算法 二、Linux线程同步 三 、条件变量 1、条件变量基本原理 2、条件变量的使用 3、条件变量使用示例 为什么 pthread_cond_wait 需要互斥量? 一、什么是死锁 死锁是计算机科学中的一个概念&#xff0c;…

Python-VBA函数之旅-type函数

目录 一、type函数的常见应用场景 二、type函数使用注意事项 三、如何用好type函数&#xff1f; 1、type函数&#xff1a; 1-1、Python&#xff1a; 1-2、VBA&#xff1a; 2、推荐阅读&#xff1a; 个人主页&#xff1a; https://myelsa1024.blog.csdn.net/ 一、type函…

设计一个游戏的基本博弈框架

设计一个游戏的基本博弈框架&#xff0c;玩家通过操作改变某个数值&#xff0c;这个数值的变动会引发一系列实时变化&#xff0c;并且当这些数值累计到特定阈值时&#xff0c;会导致游戏中出现其他变化&#xff0c;可以分为以下几个步骤&#xff1a; 1. 确定游戏类型和主题 首…

UE4_照亮环境_不同雾效的动态切换

一、问题及思路&#xff1a; 我们在一个地图上&#xff0c;经常切换不同的区域&#xff0c;不同的区域可能需要不同的色调&#xff0c;例如暖色调的野外或者幽暗的山洞&#xff0c;这两种环境上&#xff0c;雾效的选用肯定不一样&#xff0c;夕阳西下的户外用的就是偏暖的色调&…

2023年数维杯国际大学生数学建模挑战赛A题复合直升机的建模与优化控制问题解题全过程论文及程序

2023年数维杯国际大学生数学建模挑战赛 A题 复合直升机的建模与优化控制问题 原题再现&#xff1a; 直升机具有垂直起降等飞行能力&#xff0c;广泛应用于侦察、运输等领域。传统直升机的配置导致旋翼叶片在高速飞行过程中受到冲击波的影响&#xff0c;难以稳定飞行。为了在保…

558、Vue 3 学习笔记 -【常用Composition API(七)】 2024.05.13

目录 一、Composition API的优势1. Options API存在的问题2. Composition API的优势 二、 新的组件1. Fragment2. Teleport3. Suspense 三、其他1. 全局API的转移2. 其他改变 四、参考链接 一、Composition API的优势 1. Options API存在的问题 使用传统OptionsAPI中&#xf…

Rust的协程机制:原理与简单示例

在现代编程中&#xff0c;协程&#xff08;Coroutine&#xff09;已经成为实现高效并发的重要工具。Rust&#xff0c;作为一种内存安全的系统编程语言&#xff0c;也采用了协程作为其并发模型的一部分。本文将深入探讨Rust协程机制的实现原理&#xff0c;并通过一个简单的示例来…

C++|内存管理(1)

目录 C/C内存分布 堆区 栈区 静态存储区 代码区 总结 C语言中动态内存管理方式&#xff1a;malloc/calloc/realloc/free C内存管理方式 new/delete操作内置类型 new和delete操作自定义类型 operator new与operator delete函数&#xff08;重要点进行讲解&#xff09;…

R语言手把手教你进行支持向量机分析

1995年VAPINK 等人在统计学习理论的基础上提出了一种模式识别的新方法—支持向量机 。它根据有限的样本信息在模型的复杂性和学习能力之间寻求一种最佳折衷。 以期获得最好的泛化能力.支持向量机的理论基础决定了它最终求得的是全局最优值而不是局部极小值,从而也保证了它对未知…

4.2 试编写一程序,要求比较两个字符串STRING1和STRING2所含字符是否相同,若相同则显示“MATCH”,若不相同则显示“NO MATCH”

方法一&#xff1a;在程序内部设置两个字符串内容&#xff0c;终端返回是否匹配 运行效果&#xff1a; 思路&#xff1a; 1、先比较两个字符串的长度&#xff0c;如果长度不一样&#xff0c;则两组字符串肯定不匹配&#xff1b;如果长度一样&#xff0c;再进行内容的匹配 2、如…

大模型崛起与就业危机

大模型&#xff0c;特别是像我这样的人工智能&#xff0c;最有可能首先替代那些重复性高、标准化程度高、不需要太多人类直觉和情感判断的工作。这些工作通常包括数据输入、初级数据分析和处理、简单的客户服务任务等。例如&#xff0c;可以自动化的一些岗位包括&#xff1a; 1…

zabbix监控mariadb

zabbix 服务端安装请参阅&#xff1a;红帽 9 zabbix 安装流程_红帽安装zabbix-CSDN博客 源码包安装mariadb请参阅&#xff1a;源码包安装mariadb_mariadb 11 源码编译安装-CSDN博客 在MariaDB中&#xff0c;你需要创建一个专门的用户&#xff0c;用于Zabbix进行监控。这个用户…

研究幽灵漏洞及其变种(包括但不限于V1-V5)的攻击原理和基于Github的尝试

一、研究幽灵漏洞及其变种(包括但不限于V1-V5)的攻击原理 1.1 基本漏洞原理(V1) 幽灵漏洞的基本原理是由于glibc库中的gethostbyname()函数在处理域名解析时,调用了__nss_hostname_digits_dots()函数存在缓冲区溢出漏洞。 具体来说,__nss_hostname_digits_dots()使用一个固定…

绝地求生:艾伦格回归活动来了,持续近1个月,新版本皮肤、G币等奖励白嫖

嗨&#xff0c;我是闲游盒~ 29.2版本更新在即&#xff0c;新活动来啦&#xff01;目前这个活动国内官方还没发&#xff0c;我就去台湾官方搬来了中文版方便大家观看&#xff0c;也分析一下这些奖励应该怎样才能获得。 新版本将在周二进行约9小时的停机维护&#xff0c;请注意安…

JSON在线解析及格式化验证 - JSON.cn网站

JSON在线解析及格式化验证 - JSON.cn https://www.json.cn/

anaconda虚拟环境pytorch安装

1.先创建conda的虚拟环境 conda create -n gputorch python3.102.激活刚刚创建好的虚拟环境 conda activate gputorch3.设置国内镜像源 修改anaconda的源&#xff0c;即修改.condarc配置文件 .condarc在 home/用户/user/ conda config --add channels https://mirrors.tuna.…

【专利】一种日志快速分析方法、设备、存储介质

公开号CN116560938A申请号CN202310311478.5申请日2023.03.28 是我在超音速人工智能科技股份有限公司(833753) 职务作品&#xff0c;第一发明人是董事长夫妇&#xff0c;第二发明人是我。 ** 注意** &#xff1a; 内容比较多&#xff0c;还有流程图、界面等。请到 专利指定页面…

初始Django

初始Django 一、Django的历史 ​ Django 是从真实世界的应用中成长起来的&#xff0c;它是由堪萨斯&#xff08;Kansas&#xff09;州 Lawrence 城中的一个网络开发小组编写的。它诞生于 2003 年秋天&#xff0c;那时 Lawrence Journal-World 报纸的程序员 Adrian Holovaty 和…

自作聪明的AI? —— 信息处理和传递误区

一、背景 在人与人的信息传递中有一个重要问题——由于传递人主观处理不当&#xff0c;导致信息失真或产生误导。在沟通交流中&#xff0c;确实存在“自作聪明”的现象&#xff0c;即传递人在转述或解释信息时&#xff0c;根据自己对信息的理解、经验以及个人意图进行了过多的…