CUDA 两个向量相加： 当数组的元素个数大于所开的线程数量的时候的解决办法

当我们对两个向量 a
和b
相加， 结果保存到C
数组处。 传统的C代码如下：

void add(int *a, int *b, int *c) {
    int tid = 0; // this is CPU 0, so we start at zero
    while(tid < N) {
        c[tid] = a[tid] + b[tid];
        tid++;
    }
}

上述代码是在CPU上执行。 当然， 只有一个CPU， 这个CPU一次只能错做数组的一个下标的值相加。

试想， 如果假如我们的系统是multiCPUs 或者多个CPU核的时候， 具体的， 我们假设我们有两个核。 现在我们想利用好这个资源， 对其进行向量的相加。 此时我们需要写一个并行的程序。 我们可以让一个核执行数组偶数索引的相加， 从tid = 0 开始， 另一个核对数组奇数索引相加， 从tid = 1开始。 那么我们可能有如下两份代码：

CPU core 1执行：

void add(int *a, int *b, int *c) {
    int tid = 0; // this is CPU 0, so we start at zero
    while(tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += 2;
    }
}

CPU core 2 执行：

void add(int *a, int *b, int *c) {
    int tid = 1; // this is CPU 0, so we start at zero
    while(tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += 2;
    }
}

正常情况下， 我们开的线程数都是大于数组的大小的数目的， GPU程序如下：

__global__ void add(int *dev_a, int *dev_b, int *dev_c) {
    int tid = blockIdx.x; // handle the data at this index
    if(tid < N) {
        c[tid] = a[tid] + b[tid];
    }
}

由于一个GPU所能启动的线程数目是一定的。 所以加入我们的数组数目远远大于我们启动的线程的数目， 我们需要执行更改GPU的代码。 如下：

__global__ void add(int *dev_a, int *dev_b, int *dev_c) {
    int tid = blockIdx.x; // handle the data at this index
    if(tid < N) {
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x * gridDim.x;
    }
}

分析：

tid += blockDim.x * gridDim.x ， 其中blockDim.x * gridDim.x就是我们开的线程的总的数目。 这种情况下， 一般是我们开的线程数目小于数组中元素的个数N. 注意， 一个线程只能做一个数组元素的相加。 一旦线程0 完成了对 a[0]和b[0]的相加， 我们不想让他闲下来， 而是让让这个线程去做0 + blockDim.x * gridDim.x 的相加。 一次类推， 线程1 。。。, 直至这N个元素全部相加完成。

不难看出， tid + 开的总的线程数 能够解决开的线程数目小于work数目的难题。 那就是前面的线程在做完自己的work后， 去做其他的work了。