【cuda学习日记】3.2 CUDA执行模型--并行归约问题

#3.2.1 资源分配
网格和线程块大小的准则使用这些准则可以使应用程序适用于当前和将来的设备：
·保持每个块中线程数量是线程束大小（32）的倍数
·避免块太小：每个块至少要有128或256个线程
·根据内核资源的需求调整块大小
·块的数量要远远多于SM的数量，从而在设备中可以显示有足够的并行
·通过实验得到最佳执行配置和资源使用情况

串行的方法实现N个元素整数数组求和，需要N-1次求和， 需要N-1步

int sum = 0;
for (int i = 0 ; i < N; i ++){
	sum += array[i];
} 

并行加法运算：
.将输入向量划分到更小的数据块中。
2.用一个线程计算一个数据块的部分和。
3.对每个数据块的部分和再求和得出最终结果。
常用方法是使用迭代成对实现。一个数据块只包含一对元素。这种实现方式需要N-1次求和，进行log2 N步. 例子：
step1: 1+2，3+4, 5+6, 7+8;
step2: 3+7 , 11 + 15;
step3: 10 + 26
__syncthreads语句可以保证，线程块中的任一线程在进入下一次迭代之前，在当前迭代里每个线程的所有部分和都被保存在了全局内存中

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <string.h>
#include <windows.h>
#include "../common/common.h"


int recursiveReduce(int *data, int const size){
    if (size == 1) return data[0];
    int const stride = size /2;
    for (int i = 0; i < stride; i ++){
        data[i] += data[i + stride];
    }
    return recursiveReduce( data, stride);
}
__global__ void warmup( int *g_idata, int *g_odata, unsigned int n){
    unsigned int tid = threadIdx.x;

    int *idata = g_idata + blockIdx.x * blockDim.x;
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx >= n) return;  // boundary check

    for(int stride=1;stride<blockDim.x;stride*=2){
        if ((tid % (2 * stride)) == 0){
            idata[tid