霍夫变换和 OpenCL

Question

我正在尝试在 OpenCL 中为圆圈实现霍夫变换，但我遇到了非常奇怪的问题。每次我运行 Hough 内核时，我最终都会得到略微不同的累加器，即使参数相同并且累加器总是新置零的 table（例如 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);
}

这可能是我递增累加器的方式的问题吗？

Answer 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频率，因为它可能有缺陷或超频的副作用。

Answer 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)

至于累加器问题，首先我将 post 我用来创建 clBuffer 的代码（我使用 OpenCL.net library for 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");

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

上述错误导致累加器在每次执行内核时看起来都不一样并且有点畸形。

霍夫变换和 OpenCL

Hough transform and OpenCL

gpgpu

opencl

hough-transform