OpenCL - 内核运行缓慢答案

【问题标题】：OpenCL - kernel running slowOpenCL - 内核运行缓慢
【发布时间】：2013-09-14 05:10:49
【问题描述】：

我的内核非常简单。它会尝试查看代码是否有效，然后根据前缀扫描输出仅存储唯一代码：

__kernel void moveValid(__global int* sortCode, __global int* mark, __global int* processorOffsets, __global int* uniqueCode,__global int* numPoints, __global int* pointIndex) 
{
    int ig = get_global_id(0);
    int m = mark[ig];
    int j= processorOffsets[ig];
    atomic_inc(&numPoints[j-1]);
    // select
    if(m == true)
    {   
            uniqueCode[j] = sortCode[ig];
            pointIndex[j] = ig;
    }

    barrier(CLK_GLOBAL_MEM_FENCE);
}

似乎内核真的很慢。是因为if语句吗？任何人都可以就如何改进内核提供任何提示吗？这个场景也可以用select吗？

【问题讨论】：

标签： select kernel opencl

【解决方案1】：

所以，不用太深入研究您的代码，我可以就其速度给出以下反馈。我假设您使用 GPU 作为您的设备。如果您使用 CPU 作为您的设备，部分信息可能仍然适用。

atomic_inc(&numPoints[j-1]);

在大多数设备上，原子增量对于全局内存来说非常缓慢。这是因为数据必须提交到全局内存中（不能在本地缓存）。

屏障（CLK_GLOBAL_MEM_FENCE）；

此屏障确保工作组中的所有工作项在继续执行之前都存在。为什么你的代码需要这个？尤其是当无事可做时，没有理由不让您的线程完成执行。这也是一个很大的性能冲击。

如果（m == 真）

这实际上并不是我见过的最糟糕的 if 语句，因为它只有一个分支。这仍然会减慢您的代码速度，但不会像其他事情那样严重。在继续之前，将为所有串行线程（对于某些架构）计算条件。

总体而言，在此代码中，您正在执行 4 次全局内存访问，就像原子操作一样，没有数学操作。 GPU 是执行此类算法的最差设备类型，因为内存访问对全局内存非常慢，尤其是在访问被合并的情况下。您是否可以考虑将一些数组移至本地内存？

【讨论】：

谢谢。我仍然不太确定将障碍放在内核中的哪个位置以及何时。我已将 atomic_inc 更改为本地内存。内核运行得更快。
通常你不需要屏障。仅当您在工作然后想要等待您的 work_group 然后做更多工作时才需要障碍。此外，障碍可能不会像您认为的那样起作用。如果你想要一个全局屏障（例如，每个工作组完成的工作量相同），那么通过内核执行而不是使用屏障来完成。还要确保您了解不同内存（全局、本地、私有）之间的差异，而不仅仅是使用它们。