过去的一周时间研究了一下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计算优势有：
- 并行性。源自图形处理需要对多条绘制流水线的支持，所以GPU拥有了非常大的并行处理能力；
- 高密集计算。高频率、大位宽的显存 + 高速PCIe总线；
- 长流水线。一般来说目前显卡的流水线都有数百个指令阶段（而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对纹理进行渲染。也就是说通过渲染一个带有纹理的四边形，我们便可以触发着色器进行运算。最后把结果存入目标纹理中。
所以整个的过程如下：
- 把数据通过CPU组织绘制到纹理上（CPU运算）
- 使用shaders处理纹理中的数据（GPU运算）
- 进行渲染，结果绘制到目标纹理（GPU运算）
- 获得目标纹理数据，CPU处理得到结果（CPU运算）
BTW，前一段时间同事老李做的骨骼动画的VTF优化（Vertex Texture Fetch）就是采用此思路。
有了上面的这些准备后，再看CUDA和OpenCL就会容易理解很多。接下来说说这些。