Reduction并行分析:
每个线程是基于一个树状的访问模型，从上至下，上一层读取数据相加得到下一层的数据．不停的迭代，直到访问完所有的数据．
利用这么多的线程块(thread block)我们需要做的事情如下：
１． 处理非常大的数组
２． 让GPU的每个处理器保持忙碌
３． 每个thread block迭代减少数组的区域． 比如这个图，第一次是８个数据，第二次是４个．
但是碰到一个问题，在thread block中的线程可以利用同步，但是每个thread block都处理完了，CUDA中并不能提供block级别的同步机制．为什么CUDA不支持全局同步呢?由两个原因：
1. 打造高性能GPU处理器的硬件个数是非常昂贵的，处理器越多越贵．
2. 这就强制程序员尽可能少的使用block个数以避免产生死锁，（此处还为弄明白：block个数不能大于处理器个数* ）
这个问题该怎么处理呢，全局同步问题？
利用多个kernel来解决这个问题：
cuda kernel lanuch可以当做全局同步点．
cuda kernel lanuch硬件方面的消耗几乎可以忽略，软件消耗非常底.
Level0 是第一个kernel，level1 是第二个kernel．
我们的优化目标是？
1. 努力达到GPU性能极限．
2. 选择合适的度量，有两种：
GFLOP/s: (FLOPS是Floating-point Operations Per Second每秒所执行的浮点运算次数的英文缩写)用于分析计算kernel的计算性能．
Bandwidth:用于分析kernel的内存使用情况.
3. reduction是算数密集度非常低的,每个元素一个FLOP.所以我们需要优化极限带宽来提高信能.
4.以Nvida G80型号的GPU为例：
．384bit 存储接口宽度，900MHZ DDR. 384*1800(DDR 是doubel rate)/8 = 86.4GB/s
Reduction1: Interleaved Addressing
kernel代码：
__global__ void reduce0(int *g_idata, int *g_odata) { extern __shared__ int sdata[]; // each thread loads one element from global to shared mem unsigned int tid = threadIdx.x; unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; sdata[tid] = g_idata[i]; __syncthreads(); // do reduction in shared mem for(unsigned int s=1; s < blockDim.x; s *= 2) { if (tid % (2*s) == 0) { sdata[tid] += sdata[tid + s]; } __syncthreads(); } // write result for this block to global mem if (tid == 0) g_odata[blockIdx.x] = sdata[0]; }