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