假如想对输入数组执行一个简单的并行操作，并为每个输入数组元素分配一个线程，但输入数组的元素数量却超过了可用线程数，该如何处理？

如下是一个简单的 CUDA 内核，接受两个大小为 n 的输入数组，将它们相加并生成一个输出数组。

__global__ void add(int *output, const int *a, const int *b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    if (i < n){
        output[i] = a[i] + b[i];
    }
}

但是，该内核无法处理 n 大于线程数的情况。

__global__ void add(int *output, const int *a, const int *b, int n) {
    for (
        int i = blockIdx.x * blockDim.x + threadIdx.x; 
        i < n; 
        i += blockDim.x * gridDim.x
    ) {
        output[i] = a[i] + b[i];
    }
}

该内核不会假设 n 大于可用线程数，而是每次循环遍历一个网格大小的数据数组，循环的步长是网格中的线程总数（即 blockDim.x * gridDim.x ）。如果网格中有 1280 个线程，那么线程 0 将计算元素 0、1280、2560，以此类推。通过使用步长等于网格大小的循环，我们可以确保所有线程束内的寻址都是单位步长，从而获得最大的内存合并效率 。

当可用线程数大于 n 时，该内核的指令成本应该和第一个相同，因为只循环一次。

使用网格步长循环有如下好处：

• 可扩展性：可以支持任何规模的问题。
• 线程重用：可以限制使用的块数来优化性能，限制网格中的块数量后，线程会被重用以执行多次计算。线程重用可以分摊线程创建和销毁的成本，以及内核在循环前后可能执行的任何其他处理（例如线程私有或共享数据的初始化）。

参考

CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops

CUDA Grid-Stride Loops: What if you Have More Data Than Threads?