opencv - OpenCL atom_inc 分离有什么好主意吗？

Question

我想使用 OpenCL 计算图像中非零点的总数。

由于是加法工作，所以我用的是atom_inc。

这里显示了内核代码。

__kernel void points_count(__global unsigned char* image_data, __global int* total_number, __global int image_width)
{
    size_t gidx = get_global_id(0);
    size_t gidy = get_global_id(1);
    if(0!=*(image_data+gidy*image_width+gidx))
    {
        atom_inc(total_number);
    }
}

我的问题是，通过使用 atom_inc，它会变得多余吗？

每当我们遇到非零点时，我们应该等待 atom_inc。

我有这样一个想法，我们可以把整行分成几百组，我们在不同的组中找到数字，最后将它们相加。

如果我们可以这样做:

 __kernel void points_count(__global unsigned char* image_data, __global int* total_number_array, __global int image_width)
{
    size_t gidx = get_global_id(0);
    size_t gidy = get_global_id(1);
    if(0!=*(image_data+gidy*image_width+gidx))
    {
        int stepy=gidy%10;
        atom_inc(total_number_array+stepy);
    }    
}

我们会将整个问题分成更多的组。那样的话，我们就可以将total_number_array中的数字一个一个地相加。

从理论上讲，它会有很大的性能提升吧？

那么，有人对这里的求和问题有什么建议吗？

谢谢!

Answer 1

如评论中所述，这是一个缩减问题。

想法是保持单独的计数，然后在最后将它们放回一起。

考虑使用本地内存来存储值。

1. 声明一个本地缓冲区供每个工作组使用。
2. 使用 local_id 作为索引来跟踪此缓冲区中出现的次数。
3. 在执行结束时对这些值求和。