OpenCL 图像直方图答案

【问题标题】：OpenCL image histogramOpenCL 图像直方图
【发布时间】：2011-08-17 08:45:01
【问题描述】：

我正在尝试在 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 像素）上运行它时，当我总结每个通道的直方图值时，我倾向于只看到大约 1,870,000 个被记录的总值。每次也是不同的数字。我确实预料到了这一点，因为有时两个内核可能会从输出数组中获取相同的值并对其进行递增，从而有效地取消了一次递增操作（我假设？）。

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

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

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

谢谢！

【问题讨论】：

标签： opencl gpgpu gpu

【解决方案1】：

是的，您正在同时从许多工作项写入共享内存，因此如果您不以安全的方式进行更新（或更糟？只是不要这样做，您将丢失元素）。组大小的增加实际上增加了计算设备的利用率，这反过来又增加了发生冲突的可能性。因此，您最终会丢失更多更新。

但是，您似乎混淆了同步（排序线程执行顺序）和共享内存更新（通常需要原子操作或代码同步和内存屏障，以确保内存更新是对其他同步的线程可见）。

同步+屏障对您的情况不是特别有用（并且正如您所指出的，无论如何都不能用于全局同步。原因是，2 个线程组可能永远不会同时运行，因此尝试同步它们是无意义的）。它通常在所有线程开始生成公共数据集，然后都开始以不同的访问模式使用该数据集时使用。

在您的情况下，您可以使用原子操作（例如 atom_inc，请参阅 http://www.cmsoft.com.br/index.php?option=com_content&view=category&layout=blog&id=113&Itemid=168）。但是，请注意，更新高度竞争的内存地址（例如，因为您有数千个线程试图全部写入 256 个整数）可能会导致性能下降。典型的直方图代码经过的所有箍都是为了减少对直方图数据的争用。

【讨论】：

【解决方案2】：

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

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

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

[编辑：]

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

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

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

【讨论】：

【解决方案3】：

你可以检查

来自AMD Accelerated Parallel Processing (APP) SDK 的直方图示例。
OpenCL 编程指南一书的第 14 章 - 图像直方图（ISBN-13：978-0-321-74964-2）。

【讨论】：

【解决方案4】：

GPU Histogram - Sample code from Apple

【讨论】：