CUDA - 这个循环在做什么？

Question

嘿，我在一个网站上看到了这个示例内核

 __global__ void loop1( int N, float alpha, float* x, float* y ) {
   int i;
   int i0 = blockIdx.x*blockDim.x + threadIdx.x;

   for(i=i0;i<N;i+=blockDim.x*gridDim.x) {
      y[i] = alpha*x[i] + y[i];
    }
}   

在C语言中计算此函数

   for(i=0;i<N;i++) {
      y[i] = alpha*x[i] + y[i];
   }

内核中的for循环是否必要？您可以只执行y[i0] = alpha*x[i0] + y[i0]并完全删除for循环。

我只是好奇它为什么存在以及它的目的是什么。这是假设内核调用类似于loop1<<<64,256>>>>，因此可能gridDim.x = 1

- user660414

---

实际上，使用这样的内核执行配置<<<64, 256>>>>，gridDim.x值为64，而不是1，因为gridDim内置的dim3变量包含网格的维度，而此网格使用一个维度有64个线程块。 - Grzegorz Szpetkowski

Answer 1

如果您的向量条目比您启动的线程数多，则需要在内核中使用for循环。如果可能的话，启动足够的线程当然更有效。

- moggi

---

啊，我明白了，所以在这种情况下，如果N大于64*256，则需要吗？非常感谢。 - user660414

Answer 2

有趣的内核。内核中的循环是必要的，因为N大于线程总数，即16,384（blockDim.x * gridDim.x），但我认为这样做不是一个好的实践方法（CUDA的整个重点是使用SIMT概念）。根据CUDA编程指南，您最多可以有65535个线程块与一个内核。此外，从计算能力2.x（费米）开始，您可以在一个块中最多拥有1024个线程（费米之前为512）。另外，如果可能的话，您可以将代码分成多个（顺序）内核。

- Grzegorz Szpetkowski

---

有时候，拥有一个易于展开的for循环（在这种情况下是正确的）比启动太多块并不断切换它们更好。对于给定的问题，可以通过经验找到最佳数量。 - Pavan Yalamanchili

@Pavan 好的，但我认为在for循环中也可以优化(全局)内存访问。在上面的例子中，y[i] = alphax[i] + y[i]; 迭代非常分散(i0, i0 + numberOfThreads, i0+2numberOfThreads, ..)。 - Grzegorz Szpetkowski

Szepetkowski，你是对的。我通常使用for循环访问连续的数据块，所以我从来没有这样想过 :) - Pavan Yalamanchili

Answer 3

虽然我们希望相信CUDA GPU具有无限的执行资源，但事实并非如此。高度优化代码的作者发现，展开for循环（通常使用固定数量的块）可以获得最佳性能。这使编码变得痛苦，但优化的CPU代码也同样令人痛苦。

顺便说一下，评论者提到这段代码可能存在合并问题，但我不明白为什么。如果基地址正确对齐（64B，因为这些是浮点数），则该代码的所有内存事务都将被合并，前提是线程/块也可被64整除。

- ArchaeaSoftware