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_traits.hpp" 47 #include "opencv2/core/cuda/vec_math.hpp" 48 #include "opencv2/core/cuda/saturate_cast.hpp" 49 #include "opencv2/core/cuda/border_interpolate.hpp" 50 51 #include "opencv2/opencv_modules.hpp" 52 53 #ifdef HAVE_OPENCV_CUDAFILTERS 54 55 namespace cv { namespace cuda { namespace device 56 { 57 namespace imgproc 58 { 59 /////////////////////////////////////////// Corner Harris ///////////////////////////////////////////////// 60 61 texture<float, cudaTextureType2D, cudaReadModeElementType> harrisDxTex(0, cudaFilterModePoint, cudaAddressModeClamp); 62 texture<float, cudaTextureType2D, cudaReadModeElementType> harrisDyTex(0, cudaFilterModePoint, cudaAddressModeClamp); 63 cornerHarris_kernel(const int block_size,const float k,PtrStepSzf dst)64 __global__ void cornerHarris_kernel(const int block_size, const float k, PtrStepSzf dst) 65 { 66 const int x = blockIdx.x * blockDim.x + threadIdx.x; 67 const int y = blockIdx.y * blockDim.y + threadIdx.y; 68 69 if (x < dst.cols && y < dst.rows) 70 { 71 float a = 0.f; 72 float b = 0.f; 73 float c = 0.f; 74 75 const int ibegin = y - (block_size / 2); 76 const int jbegin = x - (block_size / 2); 77 const int iend = ibegin + block_size; 78 const int jend = jbegin + block_size; 79 80 for (int i = ibegin; i < iend; ++i) 81 { 82 for (int j = jbegin; j < jend; ++j) 83 { 84 float dx = tex2D(harrisDxTex, j, i); 85 float dy = tex2D(harrisDyTex, j, i); 86 87 a += dx * dx; 88 b += dx * dy; 89 c += dy * dy; 90 } 91 } 92 93 dst(y, x) = a * c - b * b - k * (a + c) * (a + c); 94 } 95 } 96 97 template <typename BR, typename BC> cornerHarris_kernel(const int block_size,const float k,PtrStepSzf dst,const BR border_row,const BC border_col)98 __global__ void cornerHarris_kernel(const int block_size, const float k, PtrStepSzf dst, const BR border_row, const BC border_col) 99 { 100 const int x = blockIdx.x * blockDim.x + threadIdx.x; 101 const int y = blockIdx.y * blockDim.y + threadIdx.y; 102 103 if (x < dst.cols && y < dst.rows) 104 { 105 float a = 0.f; 106 float b = 0.f; 107 float c = 0.f; 108 109 const int ibegin = y - (block_size / 2); 110 const int jbegin = x - (block_size / 2); 111 const int iend = ibegin + block_size; 112 const int jend = jbegin + block_size; 113 114 for (int i = ibegin; i < iend; ++i) 115 { 116 const int y = border_col.idx_row(i); 117 118 for (int j = jbegin; j < jend; ++j) 119 { 120 const int x = border_row.idx_col(j); 121 122 float dx = tex2D(harrisDxTex, x, y); 123 float dy = tex2D(harrisDyTex, x, y); 124 125 a += dx * dx; 126 b += dx * dy; 127 c += dy * dy; 128 } 129 } 130 131 dst(y, x) = a * c - b * b - k * (a + c) * (a + c); 132 } 133 } 134 cornerHarris_gpu(int block_size,float k,PtrStepSzf Dx,PtrStepSzf Dy,PtrStepSzf dst,int border_type,cudaStream_t stream)135 void cornerHarris_gpu(int block_size, float k, PtrStepSzf Dx, PtrStepSzf Dy, PtrStepSzf dst, int border_type, cudaStream_t stream) 136 { 137 dim3 block(32, 8); 138 dim3 grid(divUp(Dx.cols, block.x), divUp(Dx.rows, block.y)); 139 140 bindTexture(&harrisDxTex, Dx); 141 bindTexture(&harrisDyTex, Dy); 142 143 switch (border_type) 144 { 145 case BORDER_REFLECT101: 146 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows)); 147 break; 148 149 case BORDER_REFLECT: 150 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows)); 151 break; 152 153 case BORDER_REPLICATE: 154 cornerHarris_kernel<<<grid, block, 0, stream>>>(block_size, k, dst); 155 break; 156 } 157 158 cudaSafeCall( cudaGetLastError() ); 159 160 if (stream == 0) 161 cudaSafeCall( cudaDeviceSynchronize() ); 162 } 163 164 /////////////////////////////////////////// Corner Min Eigen Val ///////////////////////////////////////////////// 165 166 texture<float, cudaTextureType2D, cudaReadModeElementType> minEigenValDxTex(0, cudaFilterModePoint, cudaAddressModeClamp); 167 texture<float, cudaTextureType2D, cudaReadModeElementType> minEigenValDyTex(0, cudaFilterModePoint, cudaAddressModeClamp); 168 cornerMinEigenVal_kernel(const int block_size,PtrStepSzf dst)169 __global__ void cornerMinEigenVal_kernel(const int block_size, PtrStepSzf dst) 170 { 171 const int x = blockIdx.x * blockDim.x + threadIdx.x; 172 const int y = blockIdx.y * blockDim.y + threadIdx.y; 173 174 if (x < dst.cols && y < dst.rows) 175 { 176 float a = 0.f; 177 float b = 0.f; 178 float c = 0.f; 179 180 const int ibegin = y - (block_size / 2); 181 const int jbegin = x - (block_size / 2); 182 const int iend = ibegin + block_size; 183 const int jend = jbegin + block_size; 184 185 for (int i = ibegin; i < iend; ++i) 186 { 187 for (int j = jbegin; j < jend; ++j) 188 { 189 float dx = tex2D(minEigenValDxTex, j, i); 190 float dy = tex2D(minEigenValDyTex, j, i); 191 192 a += dx * dx; 193 b += dx * dy; 194 c += dy * dy; 195 } 196 } 197 198 a *= 0.5f; 199 c *= 0.5f; 200 201 dst(y, x) = (a + c) - sqrtf((a - c) * (a - c) + b * b); 202 } 203 } 204 205 206 template <typename BR, typename BC> cornerMinEigenVal_kernel(const int block_size,PtrStepSzf dst,const BR border_row,const BC border_col)207 __global__ void cornerMinEigenVal_kernel(const int block_size, PtrStepSzf dst, const BR border_row, const BC border_col) 208 { 209 const int x = blockIdx.x * blockDim.x + threadIdx.x; 210 const int y = blockIdx.y * blockDim.y + threadIdx.y; 211 212 if (x < dst.cols && y < dst.rows) 213 { 214 float a = 0.f; 215 float b = 0.f; 216 float c = 0.f; 217 218 const int ibegin = y - (block_size / 2); 219 const int jbegin = x - (block_size / 2); 220 const int iend = ibegin + block_size; 221 const int jend = jbegin + block_size; 222 223 for (int i = ibegin; i < iend; ++i) 224 { 225 int y = border_col.idx_row(i); 226 227 for (int j = jbegin; j < jend; ++j) 228 { 229 int x = border_row.idx_col(j); 230 231 float dx = tex2D(minEigenValDxTex, x, y); 232 float dy = tex2D(minEigenValDyTex, x, y); 233 234 a += dx * dx; 235 b += dx * dy; 236 c += dy * dy; 237 } 238 } 239 240 a *= 0.5f; 241 c *= 0.5f; 242 243 dst(y, x) = (a + c) - sqrtf((a - c) * (a - c) + b * b); 244 } 245 } 246 cornerMinEigenVal_gpu(int block_size,PtrStepSzf Dx,PtrStepSzf Dy,PtrStepSzf dst,int border_type,cudaStream_t stream)247 void cornerMinEigenVal_gpu(int block_size, PtrStepSzf Dx, PtrStepSzf Dy, PtrStepSzf dst, int border_type, cudaStream_t stream) 248 { 249 dim3 block(32, 8); 250 dim3 grid(divUp(Dx.cols, block.x), divUp(Dx.rows, block.y)); 251 252 bindTexture(&minEigenValDxTex, Dx); 253 bindTexture(&minEigenValDyTex, Dy); 254 255 switch (border_type) 256 { 257 case BORDER_REFLECT101: 258 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect101<void>(Dx.cols), BrdColReflect101<void>(Dx.rows)); 259 break; 260 261 case BORDER_REFLECT: 262 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst, BrdRowReflect<void>(Dx.cols), BrdColReflect<void>(Dx.rows)); 263 break; 264 265 case BORDER_REPLICATE: 266 cornerMinEigenVal_kernel<<<grid, block, 0, stream>>>(block_size, dst); 267 break; 268 } 269 270 cudaSafeCall( cudaGetLastError() ); 271 272 if (stream == 0) 273 cudaSafeCall(cudaDeviceSynchronize()); 274 } 275 } 276 }}} 277 278 #endif // HAVE_OPENCV_CUDAFILTERS 279 280 #endif // CUDA_DISABLER 281