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

Question

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

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

Answer 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 高频！

Answer 2

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

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

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