Преобразование Hough и OpenCL

Question

Я пытаюсь реализовать преобразование Hough для кругов в 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); 
} 

Может ли это быть ошибкой, как я приращение аккумулятора?

Answer 1

Мне удалось решить мою проблему, найдя и исправляя три проблемы.

Прежде всего кода ядра, строка:

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"); 

исправление здесь было просто и в значительной степени самостоятельно explainatory:

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"); 

Я заполнил таблицу аккумуляторов нулями, а затем копировал их в буфер устройства каждый раз, когда выполнял ядро.

Приведенные выше ошибки привели к тому, что аккумулятор выглядел по-разному и битово-искаженный каждый раз, когда я запускал ядро.

Answer 2

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

Alse может быть другой проблемой, если вы используете несколько устройств:

get_local_id(0) + get_group_id(0) 

здесь get_group_id (0) получает идентификатор группы для каждого устройства и начинается с 0 для всех устройств, так же как get_global_id тоже начинает ноль; поэтому при использовании нескольких устройств вы должны добавить правильные смещения в инструкции «ndrange». Несмотря на то, что разные устройства могут поддерживать одни и те же требования к точности точек плавания, один из них может дать лучшую точность, чем другие, и может дать несколько иные результаты. Если это одно устройство, вам следует попытаться снизить частоты gpu, так как это может иметь дефекты или побочные эффекты разгона.