CUDA – 夏清然的日志 https://www.qingran.net Xia Qingran Geek Blog Sun, 07 Aug 2016 09:50:33 +0000 en-US hourly 1 https://wordpress.org/?v=4.6.1 112893047 使用GPU加速H.264编码分析 https://www.qingran.net/2011/07/%e4%bd%bf%e7%94%a8gpu%e5%8a%a0%e9%80%9fh-264%e7%bc%96%e7%a0%81%e5%88%86%e6%9e%90/ https://www.qingran.net/2011/07/%e4%bd%bf%e7%94%a8gpu%e5%8a%a0%e9%80%9fh-264%e7%bc%96%e7%a0%81%e5%88%86%e6%9e%90/#comments Wed, 13 Jul 2011 22:48:46 +0000 https://www.qingran.net/?p=1215 继前面的“GPGPU”和“CUDA和OpenCL”的简介后,接下来分析一个具体的使用案例:是否可以用GPU搭建一个高性能的H.264编解码服务器?

设想一个简单的需求:

  1. 把其他编码的视频转换为指定码率的H.264;
  2. 在转换过程中做一些简单的处理(例如增删水印、字幕的处理、声音的处理等);
  3. 需要封装成指定的一种container格式,比如mp4或mkv。

 

ffmpeg完成此项工作的大概过程是:

  1. 识别文件格式,打开视频文件容器,得到video_stream;
  2. 使用libavcodec把video_stream解码成原始的frame数据;
  3. 在frame的基础上做“需求2”的处理;
  4. 编码成H.264格式;
  5. 存放到指定格式的容器中(mp4或者mkv)。

“过程1”应该是一个轻量操作,且对于现有视频格式速度应该很快(播放器看片的时候不可能先等上几秒分析一把格式再开始播放吧?)。

“过程2”需要把原有的视频编码解码并转为ffmpeg内部使用的格式,速度会比上第一步慢不少,但是应该也是可以接受的。当然在码率比较高的情况下也会非常缓慢(在Pentium4时代的机器上看1080p的高清视频很卡)。

“过程3”是需要对每一帧进行处理,那么这步是可高度并行化的应该可以放到GPU进行。

“过程4”一定是最慢的一个环节:H.264需要将图像分割成很多个宏块, 然后利用视频帧图像的帧内和帧间的相关性, 采用帧内预测或帧间预测的编码模式, 对各个宏块进行压缩。然后形成帧,组成为码流。整个过程复杂,但宏块儿的处理是可以高度并行化的操作,应该可以放到GPU进行。

“过程5”和“过程1”类似。

所以如果想用GPU加速以上的过程,那么把“过程4”和“过程2”的密集计算从CPU上转到GPU上,应该是可能的发力点。

如果要设计这个系统,软件自底层到上层有几点需要考虑:

  1. 如果显卡可以指定,应该是Nvidia。在Linux下有完善的开发运行环境,且同时支持CUDA和OpenCL;
  2. 如果选定Nvidia显卡,那么根据前文的分析,使用CUDA是更好的GPU计算架构。并且应该使用最新的CUDA 4.0版本,因为在多GPU上能更方便的开发,并且GPU直接访问内存方面也做出了改进(CUDA 4.0的新特性详细见这里)。
  3. 在决定了以上基础设施后,为了提高单机的处理能力程序应该用何种架构来编写?单进程多线程还是多进程单线程?

关键在于在程序里如何更好的调用CUDA:

  1. 单进程多线程方式
    • 每个线程进行一个视频的转换,每个线程在其线程内部直接使用CUDA;
    • 有单独的CUDA线程,其他线程当需要时通知此线程进行运算;
  2. 多进程单线程
    • 每个进程进行一个视频的转换,在其进程内部各自独立的使用CUDA;
    • 有单独的GPU进程,其他进程当需要进行GPU运算的时IPC通知此进程进行;

在我看来,2-1的方式是最简单直接的,可以直接基于ffmpeg来实现,只需单把可并行化的部分从CPU移到GPU。但是CUDA 4.0是否支持这样做?这样做效率是否最高?

硬件配置上也有以下的问题:

  1. 是否需要在一个Server上配置多块儿GPU?虽说CUDA 4.0已经声称对此有很好的支持,但是其是否能利用好?另外受限于硬件瓶颈的一些限制(CPU计算能力、内存速度、总线速度),配置几块儿GPU比较合适?
  2. 如果使用ffmpeg,那么在一台拥有16核的服务器上,理论可以同时启动16个ffmpeg进程(单线程运行)。加入GPU运算后单个视频处理速度的提高是一定的,但是否还有这么大的并行能力?以及单位时间段内的处理能力是否提升?
  3. 加入GPU后硬件成本的增加是否值得?与单独使用CPU的相比单位成本是否更低?选用哪款Nvidia显卡的性价比最高?
  4. 如何看GPU的load average和top?

最后这件事情最难的一点在于:需要对CUDA、ffmpeg、H.264、硬件都有相当深入的了解才能做好这样的一个系统

]]>
https://www.qingran.net/2011/07/%e4%bd%bf%e7%94%a8gpu%e5%8a%a0%e9%80%9fh-264%e7%bc%96%e7%a0%81%e5%88%86%e6%9e%90/feed/ 2 1215
GPGPU简介 https://www.qingran.net/2011/07/gpgpu%e7%ae%80%e4%bb%8b/ Wed, 13 Jul 2011 05:09:23 +0000 https://www.qingran.net/?p=1209 过去的一周时间研究了一下GPU做通用计算以及CUDA和OpenCL,下面会分几篇文章总结最近的成果。

图形处理单元( GPU)简称显卡是现在计算机中除CPU体系之外最复杂的一个系统。近几年来随着游戏工业的大规模发展,GPU的运算性能的增长大大超过了摩尔定律。不仅仅提高了计算机图形处理的速度和质量而且给我们提供了nb的计算平台。


GPU有两个显著的特点:

1,运算单元极多带来大并行处理能力,GPU的运算单元数量远多于CPU。例如现在民用的中端显卡Geforce GTX 570有480个CUDA cores,拥有1405 GigaFLOPS单精度浮点数运算能力。

如上图所示,CPU的大量的晶体管被cache和控制电路占用(控制指令和分支预测等操作),而ALU占用的只是一小部分。与之相反的是GPU做为专用运算器其控制电路极其简单,而且对Cache需求较小,所以大部分的晶体管用于组成各类专用运算单元和长流水线。

2,GPU拥有更快速的显存、大的显存位宽和大的显存带宽,Geforce GTX 570一般拥有1GB以上GDDR5的显存,位宽为320bit,显存频率为950MHz,且拥有152GB/s 的显存带宽。

GPU和CPU的带宽对比:

基于以上的两个特点,GPU计算优势有:

  1. 并行性。源自图形处理需要对多条绘制流水线的支持,所以GPU拥有了非常大的并行处理能力;
  2. 高密集计算。高频率、大位宽的显存 + 高速PCIe总线;
  3. 长流水线。一般来说目前显卡的流水线都有数百个指令阶段(而CPU一般只有几十个),所以GPU做为流式数据并行处理机有明显的优势。

总结一句:GPU时针对向量计算优化的并行数据处理机。

GPU编程上也有一些特殊的地方,从原理上说我们不借助CUDA或者OpenCL就能直接使用DirectX或者OpenGL直接进行GPU运算,我们可以从一个图形开发者的角度去思考如何来做。

1,纹理 == 数组

一维数组是CPU最基本的数据表达形式,同时基于此进行的偏移量计算构成了CPU编程的数据表达形式。

而对于GPU,最基本的数据排列方式是二维数组,在图形开发中,数组被做做为纹理(texture)送入GPU后进行处理。在CPU中的数组索引在GPU中改为了纹理坐标,有了纹理坐标就可以访问纹理中的每个数据,同时还必须确定坐标原点的位置。

GPU内部的额运算均是浮点数为运算,同时GPU是以一个四元组为单位进行计算的,这个四元组包括RGB三原色 + Alpha通道。

 

2,shader == kernels

传统CPU指令是以顺序、循环执行为主,例如计算1M个成员的数组的平方的结果:

for (int i=0;  i<1000000; i++)
output[i] = input[i] * input[i]

这个运算有一个非常重要的特点:那就是输入和输出的每个数组元数,它们之间是相互独立的。不管是输入的数组,还是输出结果的数组,对于同一个数组内的其他各个元素是相互独立的,我们可以不按顺序从最后一个算到第一个,或在中间任意位置选一个先算,它得到的最终结果是不变的。如果我们有一个数组运算器,或者我们有1M个CPU的话,我们便可以同一时间把整个数组一次给算出来,这样就根本不需要一个外部的循环。这就是SIMD(single instruction multiple data)。

同样的计算在OpenCL里是这样编写kernels:

const char *KernelSource = "\n" \
"__kernel void square( \n" \
" __global float* input, \n" \
" __global float* output, \n" \
" const unsigned int count) \n" \
"{ \n" \
" int i = get_global_id(0); \n" \
" output[i] = input[i] * input[i]; \n" \
"} \n" \
"\n";

可以看到,在GPU里,两个数组的相乘如此简单的一句话就搞定所有。在这里运算的GPU可编程模块,叫做片段管线(fragment pipeline),它是由多个并行处理单元组成的。在硬件和驱动逻辑中,每个数据项会被自动分配到不同的渲染线管线中去处理,到底是如何分配,则是没法编程控制的。从概念观点上看,所有对每个数据顶的运算工作都是相互独立的,也就是说不同片段在通过管线被处理的过程中,是不相互影响的。片段管线就像是一个数组处理器,它有能力一次处理一张纹理大小的数据。虽然在内部运算过程中,数据会被分割开来然后分配到不同的片段处理器中去,但是我们没办法控制片段被处理的先后顺序,我们所能知道的就是“地址”,也就是保存运算最终结果的那张纹理的纹理坐标。我们可能想像为所有工作都是并行的,没有任何的数据相互依赖性。

3,渲染 == 运算
我们准备好了纹理,写好了着色器去处理纹理,接下来我们所要做的最后一件事是告诉GPU让起来使用shader对纹理进行渲染。也就是说通过渲染一个带有纹理的四边形,我们便可以触发着色器进行运算。最后把结果存入目标纹理中。

所以整个的过程如下:

  1. 把数据通过CPU组织绘制到纹理上(CPU运算)
  2. 使用shaders处理纹理中的数据(GPU运算)
  3. 进行渲染,结果绘制到目标纹理(GPU运算)
  4. 获得目标纹理数据,CPU处理得到结果(CPU运算)

BTW,前一段时间同事老李做的骨骼动画的VTF优化(Vertex Texture Fetch)就是采用此思路。

有了上面的这些准备后,再看CUDA和OpenCL就会容易理解很多。接下来说说这些

]]>
1209