在解码中，采用GPU做解码能减少对CPU性能的消耗，尤其对1080p以上的视频。而对于编码，只要不是极端的场景，其实GPU也不输CPU，实际也有测过，可以说各有优点。并且GPU的处理速度会比CPU快很多，在大量视频需要同时处理时，采用GPU进行编解码是一种比较好的方案。
除去代码本身的问题，其中GPU编解码的性能基本会受到2个方面的影响，一个是GPU中的decoder和encoder，一般GPU专门有解码器和编码器来做处理，各型号数量不等，如果满载也会影响性能。另一个就是CPU和GPU的数据拷贝性能，本章着重分析影响该性能的因素。
目录
一、介绍CPU和GPU的关系
推荐先看我上篇由一位资深测试人员写的docker性能测试引发的思考(缓存、内存、显存、PCIe)，为了了解PCIe对性能的影响，尤其是对显卡方面的PCIe查了大量资料，这里找到一些关于写PCIe分析比较多的文章，有兴趣可以看：PCIe扫盲系列博文连载目录篇
在NVIDIA的API里，可以看到一般CPU称为Host，GPU称为Device。所以也就有了htod,dtoh,dtod代表3种不同的传输方向。
我们知道CPU和GPU的内存在物理上是独立的，并通过PCI-Express总线相连。那么这个总线就有带宽限制。如何查看显卡的PCIe带宽在我上篇也讲了，PCIe带宽最高也十几个G，而GPU之间是上百个G。
对GPU来说下行带宽RX(CPUtoGPU)和上行带宽TX(GPUtoCPU)中，其中TX是一个业内人士比较关注的瓶颈。如果需要从GPU中拿出解码后的像素帧到CPU，这个帧可是很大的，如果是1080p以上更甚。
二、影响拷贝速度的原因
1、CPU
CPU空闲时利用bandwidthtest测试出TX数据，不同显卡出来的数值不同，比如我的可能是13GB/s。而如果我将CPU跑到80%，再用bandwidthtest测试之后，RX和TX都降低到了5GB/s。
2、PCIe
由于PCIe为了降低功耗，在空闲的时候会切换到Gen1 一代的速率，当需要使用时，会自动切换到更高速率。但是如果切换的时候出问题，会导致切换到Gen1之后无法再切换到高速率。这个问题虽然情况发生的概率很低，但我在其他公司的显卡上发现过，英伟达还没有。
3、带宽短板
PCIe传输要保证的这块显卡是x16还是x8之类的，在插入显卡槽的时候不要插错了，不要把x16的卡，插到了x8槽上。
确保显卡槽是PCIe几代，一般槽上会写有。
还要确认主板上PCIe是几代的，速率是多少，这个也可以看我上篇有写。
如果还有其他因素可以评论区交流
三、内存优化
1、显卡的内存概念
首先要了解显卡的内存概念，CUDA中涉及到8中存储器类型，其中device端有6种，host端有2种。
| 存储器 | 位置 | 拥有缓存 | 访问权限 | 变量生存周期 |
|---|---|---|---|---|
| register | GPU片内 | 它本身就是，速度快 | device可读写 | 与thread相同 |
| local memory | 显存 | 无，慢 | device可读写 | 与thread相同 |
| shared memory | GPU片内 | 它本身就是，快 | device可读写 | 与block相同 |
| constant memory | 显存 | 有，快 | device可读，host可读写 | 可在程序中保持 |
| texture memory | 显存 | 有，快 | device可读，host可读写 | 可在程序中保持 |
| global memory | 显存 | 无，慢 | device可读写，host可读写 | 可在程序中保持 |
| host memory | 内存 | 无 | host可读写 | 可在程序中保持 |
| pinned memory | 内存 | 无 | host可读写 | 可在程序中保持 |
| 存储类型 | 寄存器 | 共享内存 | 纹理内存 | 常量内存 | 全局内存 |
| 带宽 | ～8TB/s | ～1.5TB/s | ～200MB/s | ～200MB/s | ～200MB/s |
| 延迟 | 1个周期 | 1～32周期 | 400～600周期 | 400～600周期 | 400～600周期 |
每个线程拥有自己的register and loacal memory;
每个线程块拥有一块shared memory共享内存;
所有线程都可以访问global memory全局内存;
还有，可以被所有线程访问的只读存储器：constant memory and texture memory
在 CUDA 中，数据复制到的显卡内存的部份，称为global memory，速度是很慢的。可以考虑将其读到shared memory里，能有效加速。
2、代码分配内存方式
CPU分配内存主要有两种方式：
- 通过C标准库中的
malloc函数完成 - 调用CUDA中的
cudaMallocHost函数
cudaMallocHost函数通过页面锁定，可以提供更高的CPU和GPU传输速率
GPU的分配方式有：
- 通过CUDA的
cudaMalloc函数 cudaMallocPitch函数cudaMallocArray函数
cudaMallocPitch 和 cudaMallocArray都没有cudaMalloc来的快。
3、统一内存寻址
延伸一个知识点：
由于之前在GPU开辟内存要调用cudaMalloc函数，而在cpu上则调用其他函数如malloc,cudaMallocHost等等，再显式调用拷贝函数的方式相互拷贝，使得代码维护难度大。
CUDA 6.0以后，出现了：统一内存寻址，unified memory简称UM，可以同时被cpu和gpu访问。
现在调用cudaMallocManaged即可完成cpu和gpu的相互访问。
但是它并不会消除拷贝的工作，只是交给CUDA自己处理内存拷贝，理论上并不会提升拷贝速度。
深入可看CUDA 6中的统一内存模型
4、主机端(CPU)内存概念
对于 CUDA 架构而言，主机端的内存可分为两种：(CUDA 之 Pinned Memory)
- 1）可分页内存（Pageable Memory）。
- 2）页锁定内存（Page-locked Memory），或称固定内存（Pinned Memory）
对于页锁定内存，操作系统不会对其进行分页和交换操作，一定是存储在物理内存，不会存储在虚拟内存，因此，GPU 可直接通过 DMA 机制，在主机和 GPU 之间快速复制数据。
在GPU上分配的内存默认都是锁页内存，这只是因为GPU不支持将内存交换到磁盘上
在主机上分配锁页内存有以下两种方式：
a 使用特殊的cudaHostAlloc函数，对用的释放内存使用cudaFreeHost函数进行内存释放
b 使用常规的malloc函数，然后将其注册为(cudaHostRegister)锁页内存,注册为锁页内存只是设置一些内部标志位以确保内存不被换出，并告诉CUDA驱动程序，该内存为锁页内存，可以直接使用而不需要使用临时缓冲区
CUDA统一内存、零复制内存、锁页内存文章里提到了一些关于锁页内存的一些点，我摘抄下来：
使用锁页内存需要注意以下几点：
a 不能分配太多，当计算机的物理内存不足或内存交换过多时，不建议使用页锁定内存，太多的话会降低系统整体性能
b 锁页内存和显存之间的拷贝速度比较快，是6G/s，普通的内存和显存之间的拷贝速度是3G/s（显存之间的拷贝速度是30G/s,CPU之间的速度是10G/s）
c 使用cudaHostAlloc函数分配内存，其内的内容需要从普通内存拷贝到锁页内存中，因此会存在：这种拷贝会带来额外的CPU内存拷贝时间开销，CPU需要把数据从可分页内存拷贝到锁页，但是采用cudaHostRegister把普通内存改为锁页内存，则不会带来额外的cpu内存拷贝时间开销，因为cudaHostAlloc的做法是先分配锁页内存，这时里面是没有数据的，那么需要将一般的内存拷贝过来，而对于cudaHostRegister内存，他是之间就使用malloc分配好的，cudaHostRegister只是设置一些内部标志位以确保其不被换出，相当于只是更改了一些标志位，就不存在前面说的数据拷贝
d 在某些设备上，设备存储器和主机锁页存储器之间的数据拷贝和内核函数可以并发执行
e 在某些设备上，可以将主机的锁页内存映射到设备地址空间，减少主机和设备之间的数据拷贝，要访问数据的时候不是像上面那那样将数据拷贝过来，而是直接通过主机总线到主机上访问 ，使用cudaHostAlloc分配时传入cudaHostAllocMapped，或者使用cudaHostRegister时传入cudaHostRegisterMapped标签
f 默认情况下，锁页内存是可以缓存的。在使用cudaHostAlloc分配时传入cudaHostAllocWriteCombined标签，将其标定为写结合，这意味着该内存没有一级二级缓存，这样有利用主机写该内存，而如果主机读取的话，速度将会极其慢，所以这种情况下的内存应当只用于那些主机只写的存储器锁页内存分配的内存，也有对应的三种形式：
cudaMemcpyToSymbol可以将数据从host拷贝到global，cudaMemcpy也是从host到>global，这种情况下二个函数有什么区别吗？
和各位大佬讨论一下后，和大家分享一下~cudaMemcpyToSymbol也有将数据从host拷贝到global的功能，以前只用过这个函数拷贝constant memory。拷贝方式的不同是由目的内存申请的方式决定的。申请的是device内存，cudaMemcpyToSymbol拷贝就是从host拷贝到global。申请的是constant内存，cudaMemcpyToSymbol拷贝就是从host拷贝到constant memory。
CUDA中与内存相关的一些函数：
cudaMemcpyToSymbol(主要用于将数据从host拷贝到device的constant memory，但也可以用于将数据从host拷贝到device的global区)这里讲下个人对cudaMalloc和cudaHostAlloc的理解：
cudaHostAlloc个人认为主要是在主机和设备之间进行数据交互的时候使用效率会比cudaMalloc高，另外需要注意的是，cudaHostAlloc分配的是主机上的内存，设备通过某种方式进行访问，cudaMalloc分配的是设备上的内存，其需要把相应的内容使用cudaMemcpy将数据从主机拷贝到设备，cudaHostAlloc函数的第三个参数有三个可选枚举，分别是：
cudaHostAllocDefault：
cudaHostAllocMapped(对应到cudaHostRegister函数是cudaHostRegisterMapped)：在该标志位下，将会分配一块可映射到设备地址空间的分页锁定主机存储器。这块存储器有两个地址，一个是主机存储器上的，一个是设备存储器上的，主机指针是cudaHostAlloc返回，设备指针通过cudaHostGetDevicePointer函数检索到，可以使用这个设备指针在内核中访问这块存储器，但是也有例外：
在支持统一虚拟地址空间的设备上，主机通过cudaHostAlloc和cudaMalloc分配的内存使用，单个虚拟地址空间，指针指向哪个存储空间（主机存储器或任意一个设备存储器），可以通过cudaPointerGetAttributes确定：
当在使用统一地址空间的设备间复制存储器时，cudaMemcpy*中的cudaMemcpyKind参数没有作用，可设置成cudaMemcpyDefault
通过cudaHostAlloc分配的存储器默认在使用统一地址空间的设备间是可分享的，cudaHostAlloc返回的指针可被在这些设备上的内核直接使用，而不需要cudaHostGetDevicePointer获得设备指针cudaHostAllocWriteCombined：主要是去掉一级二级缓存，这样主机写数据的时候会很快，而主机读取的时候极慢，所以用于就是用处主机只写不读
cudaHostRegister：是得到锁页内存的另一种方式，和cudaHostAlloc的区别是，cudaHostRegister是将现有的内存标记成锁页内存
cudaMallocPitch(opencv中的GpuMat使用的该函数对图像进行内存分配)注意的是，位于同一个Block中的线程才能实现通信，不同Block中的线程不能通过共享内存、同步进行通信，而应采用原子操作或主机介入。
5、几种存储器
这里介绍flag及对应的几种存储器，有助于对flag的理解
1)、可分享存储器(portable memory)
一块分页锁定存储器可被系统中的所有设备使用，但是默认的情况下，上面说的使用分布锁定存储器的好处只有分配它时，正在使用的设备可以享有(如果可能的话，所有的设备共享同一个地址空间)。为了让所有线程可以使用分布锁定共享存储器的好处，可以在使用cudaHostAlloc()分配时传入cudaHostAllocPortable标签，或者在使用cudaHostRegister()分布锁定存储器时，传入cudaHostRegisterPortable标签。
2)、写结合存储器
默认情况下，分页锁定主机存储器是可缓存的。可以在使用cudaHostAlloc()分配时传入cudaHostAllocWriteCombined标签使其被分配为写结合的。写结合存储器没有一级和二级缓存资源，所以应用的其它部分就有更多的缓存可用。另外写结合存储器在通过PCI-e总线传输时不会被监视(snoop)，这能够获得高达40%的传输加速。 从主机读取写结合存储器极其慢，所以写结合存储器应当只用于那些主机只写的存储器。
3)、被映射存储器
在一些设备上，在使用cudaHostAlloc()分配时传入cudaHostAllocMapped标签或者在使用cudaHostRegister()分布锁定一块主机存储器时使用cudaHostRegisterMapped标签，可分配一块被映射到设备地址空间的分页锁定主机存储器。这块存储器有两个地址：一个在主机存储器上，一个在设备存储器上。主机指针是从cudaHostAlloc()或malloc()返回的，设备指针可通过cudaHostGetDevicePointer()函数检索到，可以使用这个设备指针在内核中访问这块存储器。唯一的例外是主机和设备使用统一地址空间时。
从内核中直接访问主机存储器有许多优点：
①无须在设备上分配存储器，也不用在这块存储器和主机存储器间显式传输数据;数据传输是在内核需要的时候隐式进行的。
②无须使用流(参见3.2.5.4节)重叠数据传输和内核执行;数据传输和内核执行自动重叠。
由于被映射分页锁定存储器在主机和设备间共享，应用必须使用流或事件(参见3.2.5节)来同步存储器访问以避免任何潜在的读后写，写后读，或写后写危害。
为了在给定的主机线程中能够检索到被映射分页锁定存储器的设备指针，必须在调用任何CUDA运行时函数前调用cudaSetDeviceFlags()，并传入cudaDeviceMapHost标签。否则，cudaHostGetDevicePointer()将会返回错误。
一块分页锁定存储器可同时分配为被映射的和可分享的(前面的文章中提过)，这种情况下，每个要映射这块存储器的主机线程必须调用cudaHostGetDevicePointer()检索设备指针，因为每个主机线程持有的设备指针一般不同。
如果设备不支持被映射分页锁定存储器，cudaHostGetDevicePointer()将会返回错误。应用可以检查canMapHostMemory属性应用以查询这种能力，如果支持映射分页锁定主机存储器，将会返回1。 注意：从主机和其它设备的角度看，操作被映射分页锁定存储器的原子函数不是原子的。
四、如何解决
首先要排查PCIe的速率是不是达到最高的瓶颈，如果没到则需要依照上面查一下PCIe是否又问题。
其次，可以提高拷贝效率，那就得更换拷贝机制，尽请关注下一篇。
本博客是博主个人学习时的一些记录，个别文章加入了转载的源地址还有个别文章是汇总网上多份资料所成，在这之中也必有疏漏未加标注者，如有侵权请与博主联系。
1197
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
