CUDA：分区*巨大*问题的方法？答案

【问题标题】：CUDA: Method of partitioning *huge* problems?CUDA：分区*巨大*问题的方法？
【发布时间】：2011-08-08 04:58:53
【问题描述】：

所有这些 CUDA 百灵鸟的力量都令人头晕目眩，但我一直想知道的是一维块/网格尺寸的硬限制（通常分别为 512/65535）。

在处理范围更大（大约数十亿）的问题时，是否有一种通过内核有效设置“队列”的自动化编程方式，或者是手动切片和切块的情况？

每个人都如何处理问题划分？

【问题讨论】：

我们在CUDA中不是已经默认有队列了吗？
同样的三个人不断出现在我所有的问题中！（不抱怨）我看到到处都是提到的队列，但没有白痴的例子。
内核启动是异步的，两个连续的内核启动（在同一个流上）意味着第二个将被“排队”。实际上，我使用队列作为流的别名。
但是没有“神奇”的 queue.push(kernel>>())？我对 DIY 没有意见，但重新发明轮子并不是那么可靠：D
没有。您可以根据需要排队任意数量的内核调用 - 启动操作是非阻塞的。对于多个流，流 id 成为内核启动的一部分，即。内核>>()。默认情况下，所有内容都进入流 0，但如果您使用多个流，则可以与内核执行重叠复制，并在 Fermi 上同时运行多个内核。

标签： cuda parallel-processing gpu partitioning gpgpu

【解决方案1】：

有 2 种基本的数据分区方法，以便您可以使用 CUDA 处理它：

将数据分解为连续的块，这样每个线程都在一个块上工作。
每个线程蚕食一个数据元素。当所有线程都完成后，它们会按 numberOfThreads 移动自己并再次重复。

我已经通过简单的示例 here 解释了这些技术。对于大多数任务，方法 2 通常更容易编码和使用。

【讨论】：

【解决方案2】：

如果一维网格太小，只需使用二维（或使用 CUDA 4.0 的 Fermi 上的三维）网格代替。网格和块布局中的维度实际上只是为了方便 - 它使执行空间看起来像程序员习惯使用的那种常见数据并行输入空间（矩阵、网格、体素等）。但这只是与底层简单线性编号方案的一个非常小的抽象，可以在单个内核启动中处理超过 10^12 个唯一线程 ID。

在网格中，排序主要是列，所以如果您之前遇到过一维网格问题，“唯一的一维线程索引”计算为：

unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;

其理论上限为 512 * 65535 = 33553920 个唯一线程。等价的二维网格问题只是一维情况的简单扩展

size_t tidx = threadIdx.x + blockIdx.x * blockDim.x;
size_t tid = tidx + blockIdx.y * blockDim.x * GridDim.x;

其理论上限为 512 * 65535 * 65535 = 2198956147200 个唯一线程。 Fermi 将允许您向网格添加第三个维度，最大尺寸也为 65535，这在单个执行网格中最多提供大约 10^17 个线程。这是相当多的。

【讨论】：

我的印象是最大网格尺寸实际上就是这样；假设其他维度是单一的，则该维度的最大值。感谢您（一如既往）明确的回答，但在极端情况下 2^9*2^16*2*16 仍然不够？（考虑在 gpu 上生成、丢弃和缩减的中间数据集的流式缩减）
好吧，那么您必须设计某种分区方案来划分输入空间并按顺序处理它，或者将其分散到多个设备上。但实际上，设备内存将在可用线程数之前耗尽。当前最大的 GPU 仅附带 6Gb 内存 - 即只有 10^9 个线程，假设每个线程需要在每次内核调用的全局内存中至少对唯一的 32 位类型进行操作。