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 <thrust/device_ptr.h>
46#include <thrust/sort.h>
47
48#include "opencv2/core/cuda/common.hpp"
49#include "opencv2/core/cuda/utility.hpp"
50
51namespace cv { namespace cuda { namespace device
52{
53    namespace gfft
54    {
55        texture<float, cudaTextureType2D, cudaReadModeElementType> eigTex(0, cudaFilterModePoint, cudaAddressModeClamp);
56
57        __device__ int g_counter = 0;
58
59        template <class Mask> __global__ void findCorners(float threshold, const Mask mask, float2* corners, int max_count, int rows, int cols)
60        {
61            const int j = blockIdx.x * blockDim.x + threadIdx.x;
62            const int i = blockIdx.y * blockDim.y + threadIdx.y;
63
64            if (i > 0 && i < rows - 1 && j > 0 && j < cols - 1 && mask(i, j))
65            {
66                float val = tex2D(eigTex, j, i);
67
68                if (val > threshold)
69                {
70                    float maxVal = val;
71
72                    maxVal = ::fmax(tex2D(eigTex, j - 1, i - 1), maxVal);
73                    maxVal = ::fmax(tex2D(eigTex, j    , i - 1), maxVal);
74                    maxVal = ::fmax(tex2D(eigTex, j + 1, i - 1), maxVal);
75
76                    maxVal = ::fmax(tex2D(eigTex, j - 1, i), maxVal);
77                    maxVal = ::fmax(tex2D(eigTex, j + 1, i), maxVal);
78
79                    maxVal = ::fmax(tex2D(eigTex, j - 1, i + 1), maxVal);
80                    maxVal = ::fmax(tex2D(eigTex, j    , i + 1), maxVal);
81                    maxVal = ::fmax(tex2D(eigTex, j + 1, i + 1), maxVal);
82
83                    if (val == maxVal)
84                    {
85                        const int ind = ::atomicAdd(&g_counter, 1);
86
87                        if (ind < max_count)
88                            corners[ind] = make_float2(j, i);
89                    }
90                }
91            }
92        }
93
94        int findCorners_gpu(PtrStepSzf eig, float threshold, PtrStepSzb mask, float2* corners, int max_count)
95        {
96            void* counter_ptr;
97            cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) );
98
99            cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(int)) );
100
101            bindTexture(&eigTex, eig);
102
103            dim3 block(16, 16);
104            dim3 grid(divUp(eig.cols, block.x), divUp(eig.rows, block.y));
105
106            if (mask.data)
107                findCorners<<<grid, block>>>(threshold, SingleMask(mask), corners, max_count, eig.rows, eig.cols);
108            else
109                findCorners<<<grid, block>>>(threshold, WithOutMask(), corners, max_count, eig.rows, eig.cols);
110
111            cudaSafeCall( cudaGetLastError() );
112
113            cudaSafeCall( cudaDeviceSynchronize() );
114
115            int count;
116            cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) );
117
118            return std::min(count, max_count);
119        }
120
121        class EigGreater
122        {
123        public:
124            __device__ __forceinline__ bool operator()(float2 a, float2 b) const
125            {
126                return tex2D(eigTex, a.x, a.y) > tex2D(eigTex, b.x, b.y);
127            }
128        };
129
130
131        void sortCorners_gpu(PtrStepSzf eig, float2* corners, int count)
132        {
133            bindTexture(&eigTex, eig);
134
135            thrust::device_ptr<float2> ptr(corners);
136
137            thrust::sort(ptr, ptr + count, EigGreater());
138        }
139    } // namespace optical_flow
140}}}
141
142
143#endif /* CUDA_DISABLER */
144