OpenCL 图像直方图

Question

我正在尝试在 OpenCL 中编写直方图内核来计算 RGBA32F 输入图像的 256 个 bin R、G 和 B 直方图。我的内核看起来像这样:

const sampler_t mSampler = CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP|
                           CLK_FILTER_NEAREST;


__kernel void computeHistogram(read_only image2d_t input, __global int* rOutput,
                               __global int* gOutput, __global int* bOutput)
{

    int2 coords = {get_global_id(0), get_global_id(1)};

    float4 sample = read_imagef(input, mSampler, coords);

    uchar rbin = floor(sample.x * 255.0f);
    uchar gbin = floor(sample.y * 255.0f);
    uchar bbin = floor(sample.z * 255.0f);

    rOutput[rbin]++;
    gOutput[gbin]++;
    bOutput[bbin]++;


}

当我在 2100 x 894 图像(1,877,400 像素)上运行它时，当我对每个 channel 的直方图值求和时，我倾向于只看到大约 1,870,000 个被记录的值。每次也是不同的数字。我确实预料到了这一点，因为有时两个内核可能会从输出数组中获取相同的值并对其进行递增，从而有效地取消了一次递增操作(我假设？)。

1,870,000 输出用于 {1,1} 工作组大小(如果我没有另外指定，这似乎是默认设置的)。如果我强制使用更大的工作组大小，例如 {10,6}，我会在直方图中得到一个小得多的总和(与工作组大小的变化成正比)。这对我来说似乎很奇怪，但我猜会发生什么是组中的所有工作项同时增加输出数组值，所以它只是算作单个增量？

无论如何，我在规范中读到 OpenCL 没有全局内存同步，只有使用其 __local 内存的本地工作组内的同步。 nVidia 的直方图示例将直方图工作负载分解为一堆特定大小的子问题，计算它们的部分直方图，然后将结果合并为单个直方图。这似乎不适用于任意大小的图像。我想我可以用虚拟值填充图像数据......

作为 OpenCL 的新手，我想我想知道是否有更直接的方法可以做到这一点(因为它似乎应该是一个相对简单的 GPGPU 问题)。

谢谢!

Answer 1

如前所述，您以非同步且非原子的方式写入共享内存。这会导致错误。如果图片足够大，我有一个建议:

将您的工作组拆分为列或行的一维。使用每个内核对 col 或 row 的直方图求和，然后使用 atomic atom_inc 对其进行全局求和。这带来了私有(private)内存中最多的总结，它更快并减少了原子操作。

如果您在二维中工作，则可以在图片的某些部分上进行。

[编辑:]

我想，我有一个更好的答案:;-)

看看:http://developer.download.nvidia.com/compute/opencl/sdk/website/samples.html#oclHistogram

他们在那里有一个有趣的实现......