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像素)上运行时,当我总结每个通道的直方图值时,我倾向于只能看到记录的1,870,000个总值.每次都是不同的数字.我确实期待这一点,因为偶尔两个内核可能从输出数组中获取相同的值并递增它,有效地取消了一个增量操作(我假设？).

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

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

作为OpenCL的新手,我想我想知道是否有更简单的方法来做到这一点(因为它看起来应该是一个相对简单的GPGPU问题).

谢谢!

Answer 1

是的,你正在同时写入许多工作项的共享内存,所以如果你不以安全的方式进行更新,你会丢失元素(或者更糟糕的是？只是不要这样做).组大小的增加实际上会增加计算设备的利用率,从而增加冲突的可能性.所以你最终失去了更多的更新.

但是,您似乎混淆同步(订购线程执行顺序)和共享内存更新(通常需要原子操作,或代码同步和内存屏障),以确保内存更新对于同步的其他线程可见.

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

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

Answer 2

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

将工作组拆分为cols或rows的一维工作组.使用每个内核来总结col或行的直方图,然后使用atomic atom_inc对其进行全局求和.这使得私有内存中的总和增加得更快,并且减少了原子操作.

如果你在两个维度上工作,你可以在图片的某些部分进行.

[编辑:]

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

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

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