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