OpenCL 中的最佳本地/全局工作量

Posted 2023-02-16

【中文标题】OpenCL 中的最佳本地/全局工作量

【英文标题】：Optimal Local/Global worksizes in OpenCL

【发布时间】：2013-01-10 09:07:48

【问题描述】：

我想知道如何为 OpenCL 中的不同设备选择最佳的本地和全局工作大小？ AMD、NVIDIA、INTEL GPU 有什么通用规则吗？ 我应该分析设备的物理构造（多处理器数量、多处理器中的流处理器数量等）吗？

这是否取决于算法/实现？因为我看到一些库（如 ViennaCL）评估正确值只是测试本地/全局工作大小的许多组合并选择最佳组合。

【参考方案1】：

NVIDIA 建议您的（本地）工作组大小是 32 的倍数（等于一个 warp，这是它们的原子执行单元，这意味着 32 个线程/工作项以原子方式一起调度）。另一方面，AMD 建议使用 64 的倍数（等于一个波前）。不确定英特尔，但您可以在他们的文档中找到此类信息。

因此，当您进行一些计算并假设您有 2300 个工作项（全局大小）时，2300 不能被 64 或 32 整除。如果您不指定本地大小，OpenCL 将选择错误的本地适合你的尺寸。当您没有作为原子执行单元倍数的本地大小时会发生什么情况，您将获得空闲线程，从而导致设备利用率不佳。因此，添加一些“虚拟”线程可能是有益的，这样您就可以获得一个 32/64 的倍数的全局大小，然后使用 32/64 的局部大小（全局大小必须可以除以局部大小）。对于 2300，您可以添加 4 个虚拟线程/工作项，因为 2304 可以被 32 整除。在实际内核中，您可以编写如下内容：

int globalID = get_global_id(0);
if(globalID >= realNumberOfThreads)
globalID = 0;

这将使四个额外的线程执行与线程 0 相同的操作。（执行一些额外的工作通常比拥有许多空闲线程更快）。

希望能回答您的问题。 GL 高频！

【讨论】：

谢谢。我在 *** 的其他线程中看到了类似的建议。我想知道是否有针对 AMD 和 Intel 设备的类似建议？

您介意解释一下为什么让 4 个线程做无用的工作比让这 4 个线程空闲更快吗？

我想你误会了。假设您有 2300 个线程，并且工作组大小设置为 100，这不是 32 的倍数。因此对于每个工作组，每个工作在 GPU 上将有 4 个功能单元/线程空闲-组，共有23组。因此，您总共有 23*4 = 92 个空闲线程/功能单元。现在... 2300 个线程非常少。想象一下，当全局大小达到数百万时，您将获得多少空闲线程。

【参考方案2】：

如果您实际上是在使用少量内存进行处理（例如，存储内核私有状态），您可以为您的问题选择最直观的全局大小，并让 OpenCL 为您选择本地大小。

在这里查看我的答案：https://***.com/a/13762847/145757

如果内存管理是您算法的核心部分并且会对性能产生很大影响，您确实应该更进一步，首先检查 最大本地大小（取决于本地/私有内核的内存使用情况）使用 clGetKernelWorkGroupInfo，它本身将决定您的全局大小。