如何适当调整和启动 CUDA 网格？

Question

第一个问题：

假设我需要在计算能力为 1.3 的 Tesla C1060 上启动具有 229080 个线程的内核。

所以根据文档，这台机器有 240 个内核，每个对称多处理器上有 8 个内核，总共有 30 个 SM。

每个 SM 最多可以使用 1024 个线程，总共“并发”运行 30720 个线程。

现在，如果我定义 256 个线程的块，这意味着每个 SM 可以有 4 个块，因为 1024/256=4。因此，这 30720 个线程可以跨所有 SM 排列在 120 个块中。

现在，对于我的 229080 线程示例，我需要 229080/256=~895（向上取整）块来处理所有线程。

现在假设我想调用一个内核，我必须使用那些 229080 线程，所以我有两个选择。第一个是我划分问题，以便我在 for 循环中调用内核约 8 次，每次使用 120 个块和 30720 个线程的网格（229080/30720）。这样我可以确保设备将完全被占用。另一种选择是为整个 229080 个线程使用 895 个块的网格调用内核，在这种情况下，许多块将保持空闲状态，直到 SM 完成它拥有的 8 个块。

那么哪个是首选？这些块保持空闲等待有什么不同吗？他们会占用资源吗？

第二个问题

假设在我调用的内核中，我需要访问未合并的全局内存，因此可以选择使用共享内存。

然后我可以使用每个线程从全局内存上的数组中提取一个值，比如global_array，它的长度为 229080。现在我理解正确，你必须在复制到共享内存时避免分支，因为块上的所有线程需要拨打syncthreads() 以确保他们都可以访问共享内存。

这里的问题是，对于 229080 个线程，我需要正好 229080/256=894.84375 个块，因为有 216 个线程的剩余部分。我可以将这个数字四舍五入得到 895 个块，最后一个块将只使用 216 个线程。

但是由于我需要从长度为 229080 的 global_array 中提取值到共享内存，并且我不能使用条件语句来防止最后 40 个线程 (256-216) 访问 @987654324 上的非法地址@ 那么在使用共享内存加载时如何规避这个问题？

Answer 1

那么哪个是首选？这些块保持空闲等待有什么不同吗？他们会占用资源吗？

根据您的描述，首选单个内核。排队但未分配给 SM 的线程块不会占用您需要担心的任何资源，并且该机器绝对是为处理此类情况而设计的。 8 次内核调用的开销肯定会更慢，所有其他条件都相同。

现在我理解正确，您必须在复制到共享内存时避免分支，因为块上的所有线程都需要到达 syncthreads() 调用以确保它们都可以访问共享内存。

这个说法表面上是不正确的。您可以在复制到共享内存时进行分支。您只需要确保：

1. __syncthreads() 在分支结构之外，或者，
2. __syncthreads() 被分支构造中的所有线程访问（这实际上意味着对于块中的所有线程，the branch construct evaluates to the same path，至少在 __syncthreads() 屏障所在的点处。）

请注意，上面的选项 1 通常是可以实现的，这使得代码更易于遵循并且易于验证所有线程都可以到达屏障。

但是由于我需要从长度为 229080 的 global_array 中提取值到共享内存，并且我不能使用条件语句来防止最后 40 个线程 (256-216) 访问 global_array 上的非法地址，那么如何我在处理共享内存加载时绕过了这个问题？

做这样的事情：

int idx = threadIdx.x + (blockDim.x * blockIdx.x);
if (idx < data_size)
  shared[threadIdx.x] = global[idx];
__syncthreads();

这是完全合法的。块中的所有线程，无论是否参与将数据复制到共享内存，都会到达屏障。