0

第一个问题：

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

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

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

现在，如果我定义 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 个线程。

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

4

1 回答 1

3

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

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

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

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

  1. 是在__syncthreads()分支结构之外，或者，
  2. 分支构造中的__syncthreads()所有线程都可以到达 （这实际上意味着分支构造对于块中的所有线程评估到相同的路径，至少在__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();

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

于 2013-07-03T22:31:44.163 回答