CUDA的归约（Reduction）并行计算技术

CUDA的归约（Reduction）并行计算技术，主要用于将一组数据合并成一个单一的结果。归约操作通常涉及到对数据集的某种聚合，比如求和、求最大值、求最小值等。

归约的基本概念

• 输入: 一组数据（例如数组）。
• 输出: 一个单一的结果（例如数组的总和）。
• 操作: 通过某种操作（如加法、乘法等）将多个输入值合并为一个输出值。

CUDA中的归约

在CUDA编程中，归约操作通常在GPU上并行执行，以提高性能。其基本步骤包括：

1. 数据分配: 在GPU上分配内存并将数据复制到设备。
2. 并行处理: 使用多个线程并行处理数据，通常采用树形结构进行归约。
3. 结果收集: 将部分结果归约后，再将它们合并以得到最终结果。
4. 首先将数组分成多个部分，使用多个线程分别计算每部分的和。将部分和再进行归约，直到得到最终结果。主要的是我们这里使用的是共享内存，如果使用全局内存的话会导致访问速度较慢，可能会导致性能下降。

#include <stdio.h>
#include <cuda.h>

__global__ void reduce(int *g_input, int *g_output, int n) {
    extern __shared__ int sdata[];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    // 每个线程负责加载一个输入元素到共享内存
    if (i < n) {
        sdata[tid] = g_input[i];
    } else {
        sdata[tid] = 0; // 如果超出边界，填充0
    }
    __syncthreads(); // 确保所有线程都完成加载

    // 归约操作
    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            sdata[tid] += sdata[tid + stride];
        }
        __syncthreads(); // 确保每轮归约完成
    }

    // 仅有一个线程将结果写入输出数组
    if (tid == 0) {
        g_output[blockIdx.x] = sdata[0];
    }
}

int main() {
    const int ARRAY_SIZE = 1024;
    const int ARRAY_BYTES = ARRAY_SIZE * sizeof(int);
    int h_input[ARRAY_SIZE];
    int h_output[ARRAY_SIZE / 256]; // 假设每个块有256个线程
    int *d_input, *d_output;

    // 初始化输入数据
    for (int i = 0; i < ARRAY_SIZE; i++) {
        h_input[i] = 1; // 例如，数组初始化为1
    }

    // 分配设备内存
    cudaMalloc((void**)&d_input, ARRAY_BYTES);
    cudaMalloc((void**)&d_output, sizeof(int) * (ARRAY_SIZE / 256));

    // 复制输入数据到设备
    cudaMemcpy(d_input, h_input, ARRAY_BYTES, cudaMemcpyHostToDevice);

    // 启动归约核函数
    reduce<<<ARRAY_SIZE / 256, 256, 256 * sizeof(int)>>>(d_input, d_output, ARRAY_SIZE);

    // 复制结果回主机
    cudaMemcpy(h_output, d_output, sizeof(int) * (ARRAY_SIZE / 256), cudaMemcpyDeviceToHost);

    // 归约结果
    int final_sum = 0;
    for (int i = 0; i < ARRAY_SIZE / 256; i++) {
        final_sum += h_output[i];
    }

    printf("Sum: %d\n", final_sum);

    // 清理
    cudaFree(d_input);
    cudaFree(d_output);

    return 0;
}

g_input = [1, 2, 3, 4, 5, 6, 7, 8]
每个线程将对应的输入元素加载到共享内存 sdata 中：

sdata = [1, 2, 3, 4, 5, 6, 7, 8]

sdata = [1, 2, 3, 4, 5, 6, 7, 8]
第一次归约操作 (stride = 4)
线程 0 和 4：
sdata[0] += sdata[4]  // 1 + 5 = 6
线程 1 和 5：
sdata[1] += sdata[5]  // 2 + 6 = 8
线程 2 和 6：
sdata[2] += sdata[6]  // 3 + 7 = 10
线程 3 和 7：
sdata[3] += sdata[7]  // 4 + 8 = 12

sdata = [6, 8, 10, 12, 5, 6, 7, 8]
第二次归约操作 (stride = 2)
线程 0 和 2：
sdata[0] += sdata[2]  // 6 + 10 = 16
线程 1 和 3：
sdata[1] += sdata[3]  // 8 + 12 = 20

sdata = [16, 20, 10, 12, 5, 6, 7, 8]
第二次归约操作 (stride = 1)
线程 0 和 1：
sdata[0] += sdata[1]  // 16 + 20 = 36

sdata = [36, 20, 10, 12, 5, 6, 7, 8]

sdata[0] = 36 //等于1到8相加的和