CUDA编程练习（九） 归约

1 想法

归约算是一种比较常见的GPU上面临的问题，所以单开一个章节。

本篇以求和为归约操作。探讨的影响效率的因素如下（其实也不能叫探讨，毕竟Nsight还没弄好，因为八和九是同一天写的，服务器管理员权限还没到手，docker不行）：

1. 减少同步（书上说是减少排队，也是，单个线程中循环次数少了）；
2. 指令级并行（和1的方法论有时是一样的）；
3. 访存合并（有时和12的方法论是一样的）；
4. 原子操作（主要指全局存储的写）VS共享存储的归约；

2 代码

2.1 CPU和GPU最简单的实现

毫无解释的欲望

// GPU最简单的原子串行
__global__ void ezReduction(ELEM_TYPE* data, int n, double* result) {
    int tid = blockIdx.x*blockDim.x + threadIdx.x;

    if (tid < n) atomicAdd(result, (double)data[tid]);
}


// CPU简单的串行
double mySummation(ELEM_TYPE* vec, int n) {
    double ret=0;
    for (int i=0; i<n; i++) {
        ret += vec[i];
    }
    return ret;
}

2.2 线程内部归约

全局一次访存合并，128字节；

同时这就是指令级并行；

缓解了原子加和的问题；（但，其实，线程个数不多的时候，原子加和问题不大，但是线程要多，那么block就要多，block多了，归约才起作用。

写了两个版本，4并行和16并行。

// 指令并行，4
__global__ void reduction_ILP4(ELEM_TY