1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 //  By downloading, copying, installing or using the software you agree to this license.
6 //  If you do not agree to this license, do not download, install,
7 //  copy or use the software.
8 //
9 //
10 //                           License Agreement
11 //                For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 //   * Redistribution's of source code must retain the above copyright notice,
21 //     this list of conditions and the following disclaimer.
22 //
23 //   * Redistribution's in binary form must reproduce the above copyright notice,
24 //     this list of conditions and the following disclaimer in the documentation
25 //     and/or other materials provided with the distribution.
26 //
27 //   * The name of the copyright holders may not be used to endorse or promote products
28 //     derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42 
43 #if !defined CUDA_DISABLER
44 
45 #include "opencv2/core/cuda/common.hpp"
46 #include "opencv2/core/cuda/vec_math.hpp"
47 
48 namespace cv { namespace cuda { namespace device
49 {
50     namespace hough_segments
51     {
52         __device__ int g_counter;
53 
54         texture<uchar, cudaTextureType2D, cudaReadModeElementType> tex_mask(false, cudaFilterModePoint, cudaAddressModeClamp);
55 
houghLinesProbabilistic(const PtrStepSzi accum,int4 * out,const int maxSize,const float rho,const float theta,const int lineGap,const int lineLength,const int rows,const int cols)56         __global__ void houghLinesProbabilistic(const PtrStepSzi accum,
57                                                 int4* out, const int maxSize,
58                                                 const float rho, const float theta,
59                                                 const int lineGap, const int lineLength,
60                                                 const int rows, const int cols)
61         {
62             const int r = blockIdx.x * blockDim.x + threadIdx.x;
63             const int n = blockIdx.y * blockDim.y + threadIdx.y;
64 
65             if (r >= accum.cols - 2 || n >= accum.rows - 2)
66                 return;
67 
68             const int curVotes = accum(n + 1, r + 1);
69 
70             if (curVotes >= lineLength &&
71                 curVotes > accum(n, r) &&
72                 curVotes > accum(n, r + 1) &&
73                 curVotes > accum(n, r + 2) &&
74                 curVotes > accum(n + 1, r) &&
75                 curVotes > accum(n + 1, r + 2) &&
76                 curVotes > accum(n + 2, r) &&
77                 curVotes > accum(n + 2, r + 1) &&
78                 curVotes > accum(n + 2, r + 2))
79             {
80                 const float radius = (r - (accum.cols - 2 - 1) * 0.5f) * rho;
81                 const float angle = n * theta;
82 
83                 float cosa;
84                 float sina;
85                 sincosf(angle, &sina, &cosa);
86 
87                 float2 p0 = make_float2(cosa * radius, sina * radius);
88                 float2 dir = make_float2(-sina, cosa);
89 
90                 float2 pb[4] = {make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1), make_float2(-1, -1)};
91                 float a;
92 
93                 if (dir.x != 0)
94                 {
95                     a = -p0.x / dir.x;
96                     pb[0].x = 0;
97                     pb[0].y = p0.y + a * dir.y;
98 
99                     a = (cols - 1 - p0.x) / dir.x;
100                     pb[1].x = cols - 1;
101                     pb[1].y = p0.y + a * dir.y;
102                 }
103                 if (dir.y != 0)
104                 {
105                     a = -p0.y / dir.y;
106                     pb[2].x = p0.x + a * dir.x;
107                     pb[2].y = 0;
108 
109                     a = (rows - 1 - p0.y) / dir.y;
110                     pb[3].x = p0.x + a * dir.x;
111                     pb[3].y = rows - 1;
112                 }
113 
114                 if (pb[0].x == 0 && (pb[0].y >= 0 && pb[0].y < rows))
115                 {
116                     p0 = pb[0];
117                     if (dir.x < 0)
118                         dir = -dir;
119                 }
120                 else if (pb[1].x == cols - 1 && (pb[1].y >= 0 && pb[1].y < rows))
121                 {
122                     p0 = pb[1];
123                     if (dir.x > 0)
124                         dir = -dir;
125                 }
126                 else if (pb[2].y == 0 && (pb[2].x >= 0 && pb[2].x < cols))
127                 {
128                     p0 = pb[2];
129                     if (dir.y < 0)
130                         dir = -dir;
131                 }
132                 else if (pb[3].y == rows - 1 && (pb[3].x >= 0 && pb[3].x < cols))
133                 {
134                     p0 = pb[3];
135                     if (dir.y > 0)
136                         dir = -dir;
137                 }
138 
139                 float2 d;
140                 if (::fabsf(dir.x) > ::fabsf(dir.y))
141                 {
142                     d.x = dir.x > 0 ? 1 : -1;
143                     d.y = dir.y / ::fabsf(dir.x);
144                 }
145                 else
146                 {
147                     d.x = dir.x / ::fabsf(dir.y);
148                     d.y = dir.y > 0 ? 1 : -1;
149                 }
150 
151                 float2 line_end[2];
152                 int gap;
153                 bool inLine = false;
154 
155                 float2 p1 = p0;
156                 if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
157                     return;
158 
159                 for (;;)
160                 {
161                     if (tex2D(tex_mask, p1.x, p1.y))
162                     {
163                         gap = 0;
164 
165                         if (!inLine)
166                         {
167                             line_end[0] = p1;
168                             line_end[1] = p1;
169                             inLine = true;
170                         }
171                         else
172                         {
173                             line_end[1] = p1;
174                         }
175                     }
176                     else if (inLine)
177                     {
178                         if (++gap > lineGap)
179                         {
180                             bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
181                                              ::abs(line_end[1].y - line_end[0].y) >= lineLength;
182 
183                             if (good_line)
184                             {
185                                 const int ind = ::atomicAdd(&g_counter, 1);
186                                 if (ind < maxSize)
187                                     out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
188                             }
189 
190                             gap = 0;
191                             inLine = false;
192                         }
193                     }
194 
195                     p1 = p1 + d;
196                     if (p1.x < 0 || p1.x >= cols || p1.y < 0 || p1.y >= rows)
197                     {
198                         if (inLine)
199                         {
200                             bool good_line = ::abs(line_end[1].x - line_end[0].x) >= lineLength ||
201                                              ::abs(line_end[1].y - line_end[0].y) >= lineLength;
202 
203                             if (good_line)
204                             {
205                                 const int ind = ::atomicAdd(&g_counter, 1);
206                                 if (ind < maxSize)
207                                     out[ind] = make_int4(line_end[0].x, line_end[0].y, line_end[1].x, line_end[1].y);
208                             }
209 
210                         }
211                         break;
212                     }
213                 }
214             }
215         }
216 
houghLinesProbabilistic_gpu(PtrStepSzb mask,PtrStepSzi accum,int4 * out,int maxSize,float rho,float theta,int lineGap,int lineLength)217         int houghLinesProbabilistic_gpu(PtrStepSzb mask, PtrStepSzi accum, int4* out, int maxSize, float rho, float theta, int lineGap, int lineLength)
218         {
219             void* counterPtr;
220             cudaSafeCall( cudaGetSymbolAddress(&counterPtr, g_counter) );
221 
222             cudaSafeCall( cudaMemset(counterPtr, 0, sizeof(int)) );
223 
224             const dim3 block(32, 8);
225             const dim3 grid(divUp(accum.cols - 2, block.x), divUp(accum.rows - 2, block.y));
226 
227             bindTexture(&tex_mask, mask);
228 
229             houghLinesProbabilistic<<<grid, block>>>(accum,
230                                                      out, maxSize,
231                                                      rho, theta,
232                                                      lineGap, lineLength,
233                                                      mask.rows, mask.cols);
234             cudaSafeCall( cudaGetLastError() );
235 
236             cudaSafeCall( cudaDeviceSynchronize() );
237 
238             int totalCount;
239             cudaSafeCall( cudaMemcpy(&totalCount, counterPtr, sizeof(int), cudaMemcpyDeviceToHost) );
240 
241             totalCount = ::min(totalCount, maxSize);
242 
243             return totalCount;
244         }
245     }
246 }}}
247 
248 
249 #endif /* CUDA_DISABLER */
250