相关疑难解决方法(0)

我正在尝试在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问题).

谢谢!