GPGPU简介

July 13th, 2011 no comment

过去的一周时间研究了一下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就会容易理解很多。接下来说说这些

Comments are closed.

Leave A Comment