为何两款CUDA数组相加程序一用循环一不用却效果相同？

这两款程序之所以能达成相同的数组相加效果，核心是它们通过不同的线程-数组元素映射策略，让GPU线程完整覆盖了目标数组的所有元素，本质都是利用CUDA的并行线程层次（线程块+线程）完成计算，只是线程复用方式不同。

1. 带循环的大规模数组相加实现原理 #

这类程序的核心是单个线程复用处理多个数组元素，通常用于目标数组规模远大于GPU能启动的单批次线程数的场景：

• 程序会启动数量远小于1<<20的线程（比如仅启动512个线程），每个线程通过循环以自身线程ID + 总线程数为步长遍历数组，逐个完成元素相加。
• 示例核函数代码：

__global__ void add_with_loop(int *a, int *b, int *c, int total_elements) {
    // 计算当前线程的全局ID
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    // 循环处理多个元素，步长为总线程数
    for (int i = global_idx; i < total_elements; i += gridDim.x * blockDim.x) {
        c[i] = a[i] + b[i];
    }
}

• 假设总线程数为512，每个线程会依次处理索引为512, 1024, 1536...的元素，最终所有1<<20个元素都会被覆盖处理。

2. 无循环的小规模数组相加实现原理 #

这类程序的核心是单个线程对应单个数组元素，但要达成和前者相同的效果，实际调用时会启动足够多的线程来覆盖1<<20个元素：

• 核函数本身只处理单个元素，但启动线程时会设置线程块数×每块线程数等于1<<20（比如启动2048个线程块，每个块512个线程，总线程数为2048×512=1<<20），每个线程直接通过全局ID定位到唯一的数组元素完成相加。
• 示例核函数代码：

__global__ void add_without_loop(int *a, int *b, int *c) {
    int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[global_idx] = a[global_idx] + b[global_idx];
}

• 此时每个线程只需要处理一个元素，无需循环，最终同样覆盖所有1<<20个元素。

两者等价性的本质 #

两种实现只是线程复用策略的差异：

• 带循环的版本是「少线程复用处理多元素」，适合GPU硬件线程数有限的场景，通过循环让线程分批处理数据；
• 无循环的版本是「多线程并行处理单元素」，充分利用GPU的多线程并行能力，直接用足够多的线程一次性覆盖所有数据。
  只要两者最终处理的元素总数一致，数组相加的结果就完全相同。

内容的提问来源于stack exchange，提问作者user366312