c - cuda和矩阵运算中的Grid-Stride Loop，为什么我们需要它？

Question

__global__ void substract(float *A, float *B, float *res, int *n)
{
    int size = *n;
    int tid = threadIdx.x + blockIdx.x*blockDim.x;
    while (tid < size) 
    {
        res[tid] = A[tid] - B[tid];
        tid += blockDim.x * gridDim.x;
    }
}


int function(...) {
    int threadsPerBlock = 256;
    int blocks = (n+threadsPerBlock-1)/threadsPerBlock;
    int blocksPerGrid = 32<blocks ? 32 : blocks;
.
.
.
    substract<<<blocksPerGrid, threadsPerBlock>>>(A, B, res, n);
.
.
.
}

所以我编写了这段代码，它采用一个数组 A 来表示大小为 nxn 的矩阵和第二个数组 B 表示大小为 n 的 vector ，我将其中一个减去另一个。假设该数组的大小为 1000x1000。我是按照多个 cuda 指南上的示例编写的，但我不明白为什么我们需要这部分:tid += blockDim.x * gridDim.x;

因为它永远不会适合作为数组 id 元素，所以它总是大于或等于 1024，而我的数组只有 0-999 id，这对我来说似乎毫无用处，但如果没有它，我的程序就会崩溃，屏幕会变黑几秒钟后它返回，我收到驱动程序已恢复的弹出窗口。所以我试图理解为什么我不能只使用 tid = threadIdx.x + blockIdx.x*blockDim.x; 遍历整个数组。我在 while 循环之前打印了所有 tid，看起来它只是以随机顺序从 0 到 1024，因为它不能依赖 tid += blockDim.x * gridDim.x; 来计算我猜的数组边界内的任何内容。

Answer 1

如果您的数组大小等于或小于网格大小，那么在这种情况下网格步幅循环不会提供太多好处。当数组大小大于网格大小时，或者当您想要编写一个可以灵活处理任意数组大小而无需调整网格大小的内核时，它特别有用。

但是，如果数组大小小于或等于网格大小，则可能仍然需要网格跨度添加代码。如果您仔细考虑 while 循环，其原因就会变得显而易见:需要此加法操作来导致 while 循环在所有线程上终止。

假设您的数组大小为 1024，它等于 1024 个线程的网格大小(无论是否全部在一个 block 中；并不重要)。

最初，您的线程的 tid 索引为 0-1023。这些值都不会导致 while 循环终止。如果 while 循环永远不会终止，您的内核将挂起并永远运行，直到或除非您出现内核超时(这就是您所看到的)。

但是使用加法语句，在第一次 while 循环迭代之后，每个线程的 tid 值为 1024 或更大，这将导致所有线程的 while 循环终止(假设 >大小 为 1024 或更小)。