霍夫变换和 OpenCL答案

【问题标题】：Hough transform and OpenCL霍夫变换和 OpenCL
【发布时间】：2015-11-23 23:38:49
【问题描述】：

我正在尝试在 OpenCL 中实现圆的霍夫变换，但我遇到了非常奇怪的问题。每次我运行 Hough 内核时，我都会得到稍微不同的累加器，即使参数相同并且累加器总是一个新的归零表（例如http://imgur.com/a/VcIw1）。我的内核代码如下：

#define BLOCK_LEN 256

__kernel void HoughCirclesKernel(
    __global int* A,
    __global int* imgData,
    __global int* _width,
    __global int* _height,
    __global int* r
)
{
    __local int imgBuff[BLOCK_LEN];

    int localThreadIndex = get_local_id(0); //threadIdx.x
    int globalThreadIndex = get_local_id(0) + get_group_id(0) * BLOCK_LEN; //threadIdx.x + blockIdx.x * Block_Len
    int width = *_width; int height = *_height;
    int radius = *r;

    A[globalThreadIndex] = 0;
    barrier(CLK_GLOBAL_MEM_FENCE);

    if(globalThreadIndex < width*height)
    {
        imgBuff[localThreadIndex] = imgData[globalThreadIndex]; 
        barrier(CLK_LOCAL_MEM_FENCE);

        if(imgBuff[localThreadIndex] > 0) 
        {
            float s1, c1;
            for(int i = 0; i<180; i++)
            {
                s1 = sincos(i, &c1);
                int centerX = globalThreadIndex % width + radius * c1;
                int centerY = ((globalThreadIndex - centerX) / height) + radius * s1;

                if(centerX < width && centerY < height)
                    atomic_inc(A + centerX + centerY * width);
            }
        }
    }
    barrier(CLK_GLOBAL_MEM_FENCE);
}

这可能是我如何增加累加器的错吗？

【问题讨论】：

请发布一个可运行的示例来重现您的错误。我写一个的尝试正常运行并且每次都产生相同的结果，但是，当然，在你没有向我们展示的部分中可能会发生一些事情。顺便说一句，您可以将 height、width 和 r 作为标量传递，无需使用 1 元素数组。
P.S.如果可能，请在另一台设备上测试您的程序。此外，here's an attempt at reproduction 使用 Python 与 numpy、matplotlib 和 pyopencl。
我已经在这里上传了整个解决方案speedy.sh/m5ZXn/Hough.7z 这似乎很奇怪，因为我的另一个项目是不规则霍夫变换似乎工作得很好，圆圈是基于它的。
很遗憾，我没有VS来编译它，但是程序看起来不错。也许有 Windows 的人可以提供更多帮助。尝试的事情：在不同的设备上运行；尝试运行我的代码（上图）。检查编译器/驱动程序问题：将标量作为标量而不是数组传递；预先用零填充A，而不是在内核中；删除本地内存使用（这里不需要）；并行化累加器单元，而不是图像像素（这种方式不需要atomic_inc）。

标签： opencl gpgpu hough-transform

【解决方案1】：

if(globalThreadIndex < width*height)
{
        imgBuff[localThreadIndex] = imgData[globalThreadIndex]; 
        barrier(CLK_LOCAL_MEM_FENCE);
        ...
}

这是未定义的行为，因为分支内部存在障碍。

计算单元中的所有流单元必须进入相同的内存栅栏。

试试这个：

if(globalThreadIndex < width*height)
{
            imgBuff[localThreadIndex] = imgData[globalThreadIndex]; 

            ...
}

barrier(CLK_LOCAL_MEM_FENCE);

如果您使用多个设备，可能还会出现另一个问题：

get_local_id(0) + get_group_id(0)

这里 get_group_id(0) 获取每个设备的组 id，它从 0 开始为所有设备，就像 get_global_id 也从零开始一样；所以你应该在使用多个设备时在“ndrange”指令中添加适当的偏移量。尽管不同的设备可以支持相同的浮点精度要求，但其中一个设备可能会提供比其他设备更好的精度，并且可能会给出稍微不同的结果。如果是单机，那么你应该尝试降低gpu频率，因为它可能有缺陷或超频的副作用。

【讨论】：

【解决方案2】：

我已经找到并纠正了三个问题，从而解决了我的问题。

首先是内核代码，行：

int centerY = ((globalThreadIndex - centerX) / height) + radius * s1;

应该是：

int centerY = (globalThreadIndex / width) + radius * s1;

这里的主要变化是除以宽度，而不是高度。这导致了不准确的问题。

if(centerX < width && centerY < height)

以上条件改为：

if(x < width && x >= 0)
    if(y < height && y >=0)

至于累加器问题，我先贴出我用来创建clBuffer的代码（我用OpenCL.net library做C#）：

int[] a = new int[width*height]; //image size
ErrorCode error;
Mem cl_accumulator = (Mem)Cl.CreateBuffer(cl_context, MemFlags.ReadWrite, (IntPtr)(a.Length * sizeof(int)), out error);
CheckErr(error, "Cl.CreateBuffer");

这里的修复很简单，几乎不言自明：

int[] a = Enumerable.Repeat(0, width * height).ToArray();
ErrorCode error;
GCHandle accHandle = GCHandle.Alloc(a, GCHandleType.Pinned);
IntPtr accPtr = accHandle.AddrOfPinnedObject();
Mem cl_accumulator = (Mem)Cl.CreateBuffer(cl_context, MemFlags.ReadWrite | MemFlags.CopyHostPtr, (IntPtr)(a.Length * sizeof(int)), accPtr, out error);
CheckErr(error, "Cl.CreateBuffer");

我用零填充累加器表，然后每次执行内核时将其复制到设备缓冲区。

每次执行内核时，上述错误导致累加器看起来不同并且格式不正确。

【讨论】：