modules/imgproc/src/opencl/hough_lines.cl

   1 // This file is part of OpenCV project.
   2 // It is subject to the license terms in the LICENSE file found in the top-level directory
   3 // of this distribution and at http://opencv.org/license.html.
   4
   5 // Copyright (C) 2014, Itseez, Inc., all rights reserved.
   6 // Third party copyrights are property of their respective owners.
   7
   8 #ifdef MAKE_POINT_LIST
   9
  10 __kernel void make_point_list(__global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
  11                               __global uchar * list_ptr, int list_step, int list_offset, __global int* global_offset)
  12 {
  13     int x = get_local_id(0);
  14     int y = get_group_id(1);
  15
  16     __local int l_index;
  17     __local int l_points[LOCAL_SIZE];
  18     __global const uchar * src = src_ptr + mad24(y, src_step, src_offset);
  19     __global int * list = (__global int*)(list_ptr + list_offset);
  20
  21     if (x == 0)
  22         l_index = 0;
  23
  24     barrier(CLK_LOCAL_MEM_FENCE);
  25
  26     if (y < src_rows)
  27     {
  28         for (int i=x; i < src_cols; i+=GROUP_SIZE)
  29         {
  30             if (src[i])
  31             {
  32                 int val = (y << 16) | i;
  33                 int index = atomic_inc(&l_index);
  34                 l_points[index] = val;
  35             }
  36         }
  37     }
  38
  39     barrier(CLK_LOCAL_MEM_FENCE);
  40
  41     int offset;
  42     if (x == 0)
  43         offset = atomic_add(global_offset, l_index);
  44
  45     barrier(CLK_LOCAL_MEM_FENCE);
  46
  47     list += offset;
  48     for (int i=x; i < l_index; i+=GROUP_SIZE)
  49     {
  50         list[i] = l_points[i];
  51     }
  52 }
  53
  54 #elif defined FILL_ACCUM
  55
  56 __kernel void fill_accum(__global const uchar * list_ptr, int list_step, int list_offset,
  57                          __global uchar * accum_ptr, int accum_step, int accum_offset, int accum_rows, int accum_cols,
  58                          int count, float irho, float theta, int numrho)
  59 {
  60     int theta_idx = get_global_id(0);
  61     int count_idx = get_global_id(1);
  62     float cosVal;
  63     float sinVal = sincos(theta * theta_idx, &cosVal);
  64     sinVal *= irho;
  65     cosVal *= irho;
  66
  67     __global const int * list = (__global const int*)(list_ptr + list_offset);
  68     __global int* accum = (__global int*)(accum_ptr + mad24(theta_idx, accum_step, accum_offset));
  69     const int shift = (numrho - 1) / 2;
  70
  71     for (int i = count_idx; i < count; i += GROUP_SIZE)
  72     {
  73         const int val = list[i];
  74         const int x = (val & 0xFFFF);
  75         const int y = (val >> 16) & 0xFFFF;
  76
  77         int r = round(x * cosVal + y * sinVal) + shift;
  78         atomic_inc(accum + r + 1);
  79     }
  80 }
  81
  82 #endif
  83