0

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

__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吗？

4

1 回答 1

2

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

atomic_inc(&numPoints[j-1]);

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

屏障（CLK_GLOBAL_MEM_FENCE）；

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

如果（米 == 真）

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

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

于 2013-09-13T21:10:49.430 回答