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 // Copyright (C) 2013, OpenCV Foundation, all rights reserved.
16 // Third party copyrights are property of their respective owners.
17 //
18 // Redistribution and use in source and binary forms, with or without modification,
19 // are permitted provided that the following conditions are met:
20 //
21 //   * Redistribution's of source code must retain the above copyright notice,
22 //     this list of conditions and the following disclaimer.
23 //
24 //   * Redistribution's in binary form must reproduce the above copyright notice,
25 //     this list of conditions and the following disclaimer in the documentation
26 //     and/or other materials provided with the distribution.
27 //
28 //   * The name of the copyright holders may not be used to endorse or promote products
29 //     derived from this software without specific prior written permission.
30 //
31 // This software is provided by the copyright holders and contributors "as is" and
32 // any express or implied warranties, including, but not limited to, the implied
33 // warranties of merchantability and fitness for a particular purpose are disclaimed.
34 // In no event shall the Intel Corporation or contributors be liable for any direct,
35 // indirect, incidental, special, exemplary, or consequential damages
36 // (including, but not limited to, procurement of substitute goods or services;
37 // loss of use, data, or profits; or business interruption) however caused
38 // and on any theory of liability, whether in contract, strict liability,
39 // or tort (including negligence or otherwise) arising in any way out of
40 // the use of this software, even if advised of the possibility of such damage.
41 //
42 //M*/
43 
44 #pragma once
45 
46 #ifndef __OPENCV_CUDEV_BLOCK_VEC_DISTANCE_HPP__
47 #define __OPENCV_CUDEV_BLOCK_VEC_DISTANCE_HPP__
48 
49 #include "../common.hpp"
50 #include "../functional/functional.hpp"
51 #include "../warp/reduce.hpp"
52 #include "reduce.hpp"
53 
54 namespace cv { namespace cudev {
55 
56 //! @addtogroup cudev
57 //! @{
58 
59 // NormL1
60 
61 template <typename T> struct NormL1
62 {
63     typedef int value_type;
64     typedef uint result_type;
65 
66     result_type mySum;
67 
NormL1cv::cudev::NormL168     __device__ __forceinline__ NormL1() : mySum(0) {}
69 
reduceThreadcv::cudev::NormL170     __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
71     {
72         mySum = __sad(val1, val2, mySum);
73     }
74 
reduceWarpcv::cudev::NormL175     __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
76     {
77         warpReduce(smem, mySum, tid, plus<result_type>());
78     }
79 
reduceBlockcv::cudev::NormL180     template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
81     {
82         blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
83     }
84 
operator result_typecv::cudev::NormL185     __device__ __forceinline__ operator result_type() const
86     {
87         return mySum;
88     }
89 };
90 template <> struct NormL1<float>
91 {
92     typedef float value_type;
93     typedef float result_type;
94 
95     result_type mySum;
96 
NormL1cv::cudev::NormL197     __device__ __forceinline__ NormL1() : mySum(0.0f) {}
98 
reduceThreadcv::cudev::NormL199     __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
100     {
101         mySum += ::fabsf(val1 - val2);
102     }
103 
reduceWarpcv::cudev::NormL1104     __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
105     {
106         warpReduce(smem, mySum, tid, plus<result_type>());
107     }
108 
reduceBlockcv::cudev::NormL1109     template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
110     {
111         blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
112     }
113 
operator result_typecv::cudev::NormL1114     __device__ __forceinline__ operator result_type() const
115     {
116         return mySum;
117     }
118 };
119 
120 // NormL2
121 
122 struct NormL2
123 {
124     typedef float value_type;
125     typedef float result_type;
126 
127     result_type mySum;
128 
NormL2cv::cudev::NormL2129     __device__ __forceinline__ NormL2() : mySum(0.0f) {}
130 
reduceThreadcv::cudev::NormL2131     __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
132     {
133         const float diff = val1 - val2;
134         mySum += diff * diff;
135     }
136 
reduceWarpcv::cudev::NormL2137     __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
138     {
139         warpReduce(smem, mySum, tid, plus<result_type>());
140     }
141 
reduceBlockcv::cudev::NormL2142     template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
143     {
144         blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
145     }
146 
operator result_typecv::cudev::NormL2147     __device__ __forceinline__ operator result_type() const
148     {
149         return ::sqrtf(mySum);
150     }
151 };
152 
153 // NormHamming
154 
155 struct NormHamming
156 {
157     typedef int value_type;
158     typedef int result_type;
159 
160     result_type mySum;
161 
NormHammingcv::cudev::NormHamming162     __device__ __forceinline__ NormHamming() : mySum(0) {}
163 
reduceThreadcv::cudev::NormHamming164     __device__ __forceinline__ void reduceThread(value_type val1, value_type val2)
165     {
166         mySum += __popc(val1 ^ val2);
167     }
168 
reduceWarpcv::cudev::NormHamming169     __device__ __forceinline__ void reduceWarp(result_type* smem, uint tid)
170     {
171         warpReduce(smem, mySum, tid, plus<result_type>());
172     }
173 
reduceBlockcv::cudev::NormHamming174     template <int THREAD_DIM> __device__ __forceinline__ void reduceBlock(result_type* smem, uint tid)
175     {
176         blockReduce<THREAD_DIM>(smem, mySum, tid, plus<result_type>());
177     }
178 
operator result_typecv::cudev::NormHamming179     __device__ __forceinline__ operator result_type() const
180     {
181         return mySum;
182     }
183 };
184 
185 //! @}
186 
187 }}
188 
189 #endif
190