CUDA：我应该如何处理线程数不能表示为 dimGrid*dimBlock 的情况？答案

【问题标题】：CUDA: How should I handle cases where the number of threads cannot be represented as a dimGrid*dimBlock?CUDA：我应该如何处理线程数不能表示为 dimGrid*dimBlock 的情况？
【发布时间】：2014-03-17 07:20:09
【问题描述】：

假设我的输入包含 7 个数据点，在这些数据点上执行了一些计算，并将结果写回大小为 7 的输出数组。将块维度声明为 4 会导致网格大小为 2，这会导致尝试使用无效的线程 id（使用 pt_id=blockIdx.x*blockDim.x+threadID.x）为 7 运行内核，并且由于无效的内存访问而失败（因为我根据线程 id 访问了一些数组） . 我可以在我的内核中添加代码，专门将线程 id 与 max_thread_id 参数进行比较，如果 thread_id>max_thread_id 则不执行任何操作，但我想知道是否有更漂亮的方法来处理不规则的输入数组。

【问题讨论】：

你应该将数组大小传递给内核，比如N，并检查是否pt_id<N。我不确定你想用max_thread_id 实现什么。我认为您的帖子标题不合适。线程总数始终可以表示为dimGrid*dimBlock。这让我觉得您在数组大小和网格大小之间产生了一些混淆。
更漂亮是什么意思？普通的线程检查有什么问题？它使用得相当普遍。
杰克 - 我的 max_thread_id 正是你的 N。澄清一下 - 我的问题是关于 (dimGrid-1)*dimBlock
您可能会发现这篇文章很有帮助：CUDA Pro Tip: Write Flexible Kernels with Grid-Stride Loops

标签： cuda

【解决方案1】：

任务的大小不是块维度的倍数是很常见的事情。我最常用的解决方案是这个。假设您的输入数据大小为N，并且您希望将启动配置为等于BLOCK_SIZE 的块大小。在这种情况下，您的启动配置可能如下所示：

kernel_function<<<(N + BLOCK_SIZE - 1) / BLOCK_SIZE, BLOCK_SIZE>>>(...);

在内核代码中，每个线程决定它是否应该做一些工作，像这样：

int id = blockIdx.x*blockDim.x + threadIdx.x;
if (id < N) { /* do the stuff */ }
else { return; }

如果任务的大小 (N) 取决于输入，您也必须将此值作为参数传递给内核函数。此外，将N 和BLOCK_SIZE 的值定义为宏或模板参数是很常见的。

最后，如果您的输入数组尺寸很小，就像您的示例中那样，GPU 仍然没有得到充分利用，并行性不会给您带来任何好处，甚至会降低算法的性能。

【讨论】：

感谢 stuhlo，这是我最终实施的解决方案。 7 只是一个例子，实际输入数组的大小是 CUDA 值得的。