运行并行 OpenCL 内核答案

【问题标题】：Running parallel OpenCL kernels运行并行 OpenCL 内核
【发布时间】：2015-01-30 04:30:12
【问题描述】：

我一直在研究 OpenCL，看看它在我的上下文中是否有用，虽然我了解基础知识，但我不确定我是否了解如何强制内核的多个实例在其中运行并行。

在我的情况下，我想要运行的应用程序本质上是顺序的，并且（在某些情况下）需要非常大的输入（数百 MB）。但是，有问题的应用程序有许多可以设置的不同选项/标志，在某些情况下可以使其更快或更慢。我希望我们可以为 OpenCL 重新编写应用程序，然后并行执行每个选项/标志，而不是猜测要使用哪些标志集。

我的问题是：显卡可以并行运行多少个内核。这是购买时可以看的东西吗？它是否与着色器的数量、内存或应用程序/内核的大小有关？

此外，虽然应用程序的输入是相同的，但每次执行都会以不同的方式修改数据。我是否需要将输入数据分别传输到每个内核以允许这样做，或者每个内核都可以分配“本地”内存。

最后，这甚至需要多个内核，我可以使用工作项来代替吗？在这种情况下，您如何确定可以并行运行多少个工作项？

（参考：http://www.drdobbs.com/parallel/a-gentle-introduction-to-opencl/231002854?pgno=3）

【问题讨论】：

标签： parallel-processing opencl

【解决方案1】：

您的问题似乎不时出现在各种论坛和 SO 上。您将用于在硬件级别上单独运行内核的功能称为设备裂变。在this page 或谷歌“cl_ext_device_fission”上阅读更多关于扩展的信息。

此扩展已在 CPU 上启用了很长时间，但在 GPU 上未启用。最新的图形硬件可能支持设备裂变。您可能至少需要 2014 年第二季度或更新的 GPU，但这必须由您自己研究。

仅使用 OpenCL 软件让内核并行运行的方法是将它们与同一设备上的不同命令队列一起排队。一些开发人员说多队列会损害性能，但我个人没有经验。

【讨论】：

感谢您的回答，事实证明我根本不需要运行多个内核，只需将一些全局数据传入内核访问的所有实例，然后将自定义数据传入数组中。但是，我最感兴趣的是确定可以并行运行多少个内核实例——据我所知，这与内核使用的资源量有关。有没有办法检查这个？

【解决方案2】：

一张显卡可以并行运行多少个内核？

您可以查看可以在显卡上并行运行多少个内核实例（即具有不同启动 ID 的相同内核代码）。这是 SIMD/CU/着色器/等的功能。取决于 GPU 供应商喜欢如何称呼它们。获得真正执行的内核实例的确切数量有点复杂，因为这取决于 占用率，这取决于内核使用的资源，例如使用的寄存器，使用的本地内存。

如果您的意思是可以并行运行多少个内核调度（即不同的内核代码和 cl_kernel 对象或不同的内核参数），那么我所知道的所有 GPU 只能运行一个内核一次。这些内核可以从多个命令队列中提取，但 GPU 一次只能处理一个。这就是当前 GPU 不支持 cl_ext_device_fission 的原因 - 无法“拆分”硬件。不过，您可以在内核代码中自己完成（见下文）。

每个内核都可以分配“本地”内存吗？

是的。这正是 OpenCL 本地内存的用途。但是，它是一种有限的资源，因此应将其视为内核控制的缓存而不是堆。

在这种情况下，您如何确定可以并行运行多少工作项？

假设内核实例，与第一个问题的答案相同。

这甚至需要多个内核，我可以使用工作项来代替吗？

您可以使用超级内核来模拟不同的内核运行，该内核根据工作项全局 ID 决定运行哪个子内核。例如：

void subKernel0( .... )
{
    int gid = get_global_id(0);
    // etc.
}

void subKernel1( .... )
{
    int gid = get_global_id(0) - DISPATCH_SIZE_0;
    // etc.
}

__kernel uberKernel( .... )
{
    if( get_global_id(0) < DISPATCH_SIZE_0 )
    {
        subKernel0( .... );
    }
    else if( get_global_id(0) < DISPATCH_SIZE_0 + DISPATCH_SIZE_1 )
    {
        subKernel1( .... );
    }
    else if( .... )
    {
        // etc.
    }
}

将调度大小设为 32/64 倍数等的常用性能建议也适用于此。您还必须调整其他各种 id。

【讨论】：

根据 AMD 的说法，GCN 1.0 卡可以并行执行内核：AMD OpenCL Programming Guide, section 5.7.4:“AMD Southern Islands GPU 可以在没有依赖关系的情况下同时执行多个内核。”

【解决方案3】：

为了支持 2008 至 2015 硬件的兼容性，只需安全地假设每个 gpu 在任何时刻只能运行一个内核，并且内核在 runtume 上交换和编译，排队以模拟多个内核。内核的交换是大内核比小内核更好的原因。

单内核客户端计算单元是默认值。具有同时运行 2 个并行不同的独立内核的选项是例外。假设它很少见且不受支持或速度较慢。

当然，一台计算机中的 2cpus 可以做到这一点。但截至 2016 年，在一个系统中拥有 2 个 CPU 仍然有点太少见了。更罕见的是有 4 个。

某些显卡可能无法并行运行 2 个内核。假设他们不是这样的。

【讨论】：