GPGPU – 夏清然的日志 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
CUDA和OpenCL https://www.qingran.net/2011/07/cuda%e5%92%8copencl/ https://www.qingran.net/2011/07/cuda%e5%92%8copencl/#comments Wed, 13 Jul 2011 20:25:11 +0000 https://www.qingran.net/?p=1213 接前文“GPGPU”。

虽然我们可以使用已有的图形API来调用GPU,但是通过前文的分析,这个过程冗长且复杂。严重违反了程序员的优雅、和lazy原则。需要去学习图形学的一些知识,了解texture、shader等图形学专用概念,而且需要学习CGSL或者HLSL等shader着色语言。而且还要熟悉OpenGL和DirectX等图形学API,这一箩筐的知识没有一个一年半载是搞不定的。而且这样的方式不符合正常程序的编写习惯,所以难以优化。

最为3D程序员着想的Nvidia终于在07年搞出 CUDA ,这个架构的编程基于C 语言来进行,程序员只需要熟习CUDA的内存布局、线程层级,就能完成 CUDA 程序的开发,彻底跟那一箩筐图形学知识脱了干系,这个世界清静了。。。

CUDA 的推出是革命性的,而且OpenCL和M$的DirectX Computer GPGPU API 都基本上是采用了 CUDA 的架构和编程模型,可以说熟习 CUDA是OpenCL 的快速入门之道(Apple你肿了么?最新的mac为啥都用A卡了?)。CUDA的架构图(取自这里):

自下往上CUDA架构由四部分组成:

  1. 显卡内部的GPU的并行计算引擎;
  2. OS内核态的硬件支持,包括初始化,配置等;
  3. 用户级的驱动程序,为开发者提供的设备级的API接口;
  4. PTX指令集架构(ISA)的并行计算内核和功能.

对于开发者来说,CUDA提供了以下内容:

  • Libraries:高层次的为CUDA专门优化的BLAS, FFT等公用函数;
  • C Runtime:CUDA C语言运行库提供了支持在GPU上执行标准C语言语法喝函数,并允许其他高级语言如Fortran,Java和Python的调用;
  • Tools:NVIDIA C Compiler (nvcc),  CUDA Debugger (cudagdb),  CUDA Visual Profiler (cudaprof)

 

CUDA对OpenCL和其他编程语言的支持

由上图可知,Nvidia的OpenCL是基于CUDA Driver构建起来的,如果我们使用OpenCL那么就根据其规范编写类C语言的OpenCL内核,CUDA的Driver会接受这些计算请求并翻译给GPU运行当使用的设备级的编程接口。MS在Nvidia上的Direct Computer也是基于此,如果使用DirectX,那么需要用HLSL.编写DirectX的computer shaders,然后调CUDA Driver.

同时CUDA C Runtime的存在使我们方便的使用其他高级语言进行CUDA程序的编写,我目前常用Python就有PyCUDA的包装。

下面的内容都以直接使用CUDA的标准kernels并调用CUDA Driver API进行,这也是最直接的方法。

 

 

运算方式
CUDA支持大量的线程级并行(Thread Level Parallel),因为在GPU硬件中动态地创建、调度和执行这些线程是非常轻量的(而在CPU中,这些操作都是重量级的)。CUDA编程模型可以看作是将GPU做为协处理器的系统,CPU来控制程序整体的逻辑和调度,GPU来运行一些能够被高度线程化的数据并行部分。

CUDA程序包括串行计算部分和并行计算部分,并行计算部分称为Kernels,内核是一个可以在GPU执行的并行代码段。理想情况下,运行于CPU上的串行代码的只是清理上个内核函数,并启动下一个内核函数。

CUDA Kernels只对ANSI C进行了很小的扩展,以实现线程按照两个层次进行组织、共享存储器和栅栏同步。这些关键特性使得CUDA拥有了两个层次的并行:线程级并行实现的细粒度数据并行,和任务级并行实现的粗粒度并行。关于CUDA的线程模型需要单独拉出来说了(我看了半天文档确实没看明白,线程模型可能是CUDA最关键的一个部分,需要买块儿Geforce GTX 5xx的显卡做一个实际实验环境啊~)

 

内存操作

在 CUDA 中,GPU 不能直接存取主内存,只能存取显卡上的显存。因此,需要将数据从主内存先复制到显卡内存中,进行运算后,再将结果从显存中复制到主内存中。这些 复制的动作会受限于 PCI Express 的速度。使用 PCI Express x16 时,PCI Express 1.0 可以提供双向各 4GB/s 的带宽,而 PCI Express 2.0 则可提供 8GB/s 的带宽。当然这都是理论值。

但刚刚发布的CUDA 4.0对内存地址映射有了更好的支持:任意在CPU的malloc/new的内存,都可以通过调用cudaHostRegister注册给GPU,同样也有cudaHostUnregister来取消这个注册,而不需要cudaMemcpy。PCI-E本身就支持把一个host地址映射到device端,且DirectX和OpenGL均用这种方式来处理texture大数据。另外这个特性更方便我们编写代码:可以把一个大的CPU程序一点一点地改成GPU程序,并一步一步对比验证结果。以前这件事情需要前后都插入不少cudaMalloc和cudaMemcpy等代码,繁琐而且容易出错。

基于CUDA的复杂线程模型,内存管理的细节应该也是相当复杂,以后专文说吧。

 

执行方式

线程模型和内存管理没整懂,runtime时怎么做的就更不明白了,专文说把。

 

]]>
https://www.qingran.net/2011/07/cuda%e5%92%8copencl/feed/ 4 1213
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