Point Cloud Library (PCL)  1.14.1
NCVPixelOperations.hpp
1 /*
2  * Software License Agreement (BSD License)
3  *
4  * Point Cloud Library (PCL) - www.pointclouds.org
5  * Copyright (C) 2009-2010, NVIDIA Corporation, all rights reserved.
6  * Third party copyrights are property of their respective owners.
7  *
8  * All rights reserved.
9  *
10  * Redistribution and use in source and binary forms, with or without
11  * modification, are permitted provided that the following conditions
12  * are met:
13  *
14  * * Redistributions of source code must retain the above copyright
15  * notice, this list of conditions and the following disclaimer.
16  * * Redistributions in binary form must reproduce the above
17  * copyright notice, this list of conditions and the following
18  * disclaimer in the documentation and/or other materials provided
19  * with the distribution.
20  * * Neither the name of Willow Garage, Inc. nor the names of its
21  * contributors may be used to endorse or promote products derived
22  * from this software without specific prior written permission.
23  *
24  * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
25  * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
26  * LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
27  * FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
28  * COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
29  * INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
30  * BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
31  * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
32  * CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
33  * LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
34  * ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
35  * POSSIBILITY OF SUCH DAMAGE.
36  *
37  * $Id: $
38  * Ported to PCL by Koen Buys : Attention Work in progress!
39  */
40 
41 #ifndef _ncv_pixel_operations_hpp_
42 #define _ncv_pixel_operations_hpp_
43 
44 #include <limits>
45 #include "NCV.hpp"
46 
47 template<typename TBase> inline __host__ __device__ TBase _pixMaxVal();
48 template<> inline __host__ __device__ Ncv8u _pixMaxVal<Ncv8u>() {return std::numeric_limits<unsigned char>::max();}
49 template<> inline __host__ __device__ Ncv16u _pixMaxVal<Ncv16u>() {return std::numeric_limits<unsigned short>::max();}
50 template<> inline __host__ __device__ Ncv32u _pixMaxVal<Ncv32u>() {return std::numeric_limits<unsigned int>::max();}
51 template<> inline __host__ __device__ Ncv8s _pixMaxVal<Ncv8s>() {return std::numeric_limits<signed char>::max();}
52 template<> inline __host__ __device__ Ncv16s _pixMaxVal<Ncv16s>() {return std::numeric_limits<short>::max();}
53 template<> inline __host__ __device__ Ncv32s _pixMaxVal<Ncv32s>() {return std::numeric_limits<int>::max();}
54 template<> inline __host__ __device__ Ncv32f _pixMaxVal<Ncv32f>() {return std::numeric_limits<float>::max();}
55 template<> inline __host__ __device__ Ncv64f _pixMaxVal<Ncv64f>() {return std::numeric_limits<double>::max();}
56 
57 template<typename TBase> inline __host__ __device__ TBase _pixMinVal();
58 template<> inline __host__ __device__ Ncv8u _pixMinVal<Ncv8u>() {return 0;}
59 template<> inline __host__ __device__ Ncv16u _pixMinVal<Ncv16u>() {return 0;}
60 template<> inline __host__ __device__ Ncv32u _pixMinVal<Ncv32u>() {return 0;}
61 template<> inline __host__ __device__ Ncv8s _pixMinVal<Ncv8s>() {return std::numeric_limits<signed char>::min();}
62 template<> inline __host__ __device__ Ncv16s _pixMinVal<Ncv16s>() {return std::numeric_limits<short>::min();}
63 template<> inline __host__ __device__ Ncv32s _pixMinVal<Ncv32s>() {return std::numeric_limits<int>::min();}
64 template<> inline __host__ __device__ Ncv32f _pixMinVal<Ncv32f>() {return std::numeric_limits<float>::min();}
65 template<> inline __host__ __device__ Ncv64f _pixMinVal<Ncv64f>() {return std::numeric_limits<double>::min();}
66 
67 template<typename Tvec> struct TConvVec2Base;
68 template<> struct TConvVec2Base<uchar1> {using TBase = Ncv8u;};
69 template<> struct TConvVec2Base<uchar3> {using TBase = Ncv8u;};
70 template<> struct TConvVec2Base<uchar4> {using TBase = Ncv8u;};
71 template<> struct TConvVec2Base<ushort1> {using TBase = Ncv16u;};
72 template<> struct TConvVec2Base<ushort3> {using TBase = Ncv16u;};
73 template<> struct TConvVec2Base<ushort4> {using TBase = Ncv16u;};
74 template<> struct TConvVec2Base<uint1> {using TBase = Ncv32u;};
75 template<> struct TConvVec2Base<uint3> {using TBase = Ncv32u;};
76 template<> struct TConvVec2Base<uint4> {using TBase = Ncv32u;};
77 template<> struct TConvVec2Base<float1> {using TBase = Ncv32f;};
78 template<> struct TConvVec2Base<float3> {using TBase = Ncv32f;};
79 template<> struct TConvVec2Base<float4> {using TBase = Ncv32f;};
80 template<> struct TConvVec2Base<double1> {using TBase = Ncv64f;};
81 template<> struct TConvVec2Base<double3> {using TBase = Ncv64f;};
82 template<> struct TConvVec2Base<double4> {using TBase = Ncv64f;};
83 
84 #define NC(T) (sizeof(T) / sizeof(TConvVec2Base<T>::TBase))
85 
86 template<typename TBase, Ncv32u NC> struct TConvBase2Vec;
87 template<> struct TConvBase2Vec<Ncv8u, 1> {using TVec = uchar1;};
88 template<> struct TConvBase2Vec<Ncv8u, 3> {using TVec = uchar3;};
89 template<> struct TConvBase2Vec<Ncv8u, 4> {using TVec = uchar4;};
90 template<> struct TConvBase2Vec<Ncv16u, 1> {using TVec = ushort1;};
91 template<> struct TConvBase2Vec<Ncv16u, 3> {using TVec = ushort3;};
92 template<> struct TConvBase2Vec<Ncv16u, 4> {using TVec = ushort4;};
93 template<> struct TConvBase2Vec<Ncv32u, 1> {using TVec = uint1;};
94 template<> struct TConvBase2Vec<Ncv32u, 3> {using TVec = uint3;};
95 template<> struct TConvBase2Vec<Ncv32u, 4> {using TVec = uint4;};
96 template<> struct TConvBase2Vec<Ncv32f, 1> {using TVec = float1;};
97 template<> struct TConvBase2Vec<Ncv32f, 3> {using TVec = float3;};
98 template<> struct TConvBase2Vec<Ncv32f, 4> {using TVec = float4;};
99 template<> struct TConvBase2Vec<Ncv64f, 1> {using TVec = double1;};
100 template<> struct TConvBase2Vec<Ncv64f, 3> {using TVec = double3;};
101 template<> struct TConvBase2Vec<Ncv64f, 4> {using TVec = double4;};
102 
103 //TODO: consider using CUDA intrinsics to avoid branching
104 template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a);};
105 template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a, 0, std::numeric_limits<unsigned short>::max());}
106 template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a, 0, std::numeric_limits<unsigned int>::max());}
107 template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}
108 
109 //TODO: consider using CUDA intrinsics to avoid branching
110 template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a+0.5f);}
111 template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a+0.5f, 0, std::numeric_limits<unsigned short>::max());}
112 template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a+0.5f, 0, std::numeric_limits<unsigned int>::max());}
113 template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}
114 
115 template<typename Tout> inline Tout _pixMakeZero();
116 template<> inline __host__ __device__ uchar1 _pixMakeZero<uchar1>() {return make_uchar1(0);}
117 template<> inline __host__ __device__ uchar3 _pixMakeZero<uchar3>() {return make_uchar3(0,0,0);}
118 template<> inline __host__ __device__ uchar4 _pixMakeZero<uchar4>() {return make_uchar4(0,0,0,0);}
119 template<> inline __host__ __device__ ushort1 _pixMakeZero<ushort1>() {return make_ushort1(0);}
120 template<> inline __host__ __device__ ushort3 _pixMakeZero<ushort3>() {return make_ushort3(0,0,0);}
121 template<> inline __host__ __device__ ushort4 _pixMakeZero<ushort4>() {return make_ushort4(0,0,0,0);}
122 template<> inline __host__ __device__ uint1 _pixMakeZero<uint1>() {return make_uint1(0);}
123 template<> inline __host__ __device__ uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,0);}
124 template<> inline __host__ __device__ uint4 _pixMakeZero<uint4>() {return make_uint4(0,0,0,0);}
125 template<> inline __host__ __device__ float1 _pixMakeZero<float1>() {return make_float1(0.f);}
126 template<> inline __host__ __device__ float3 _pixMakeZero<float3>() {return make_float3(0.f,0.f,0.f);}
127 template<> inline __host__ __device__ float4 _pixMakeZero<float4>() {return make_float4(0.f,0.f,0.f,0.f);}
128 template<> inline __host__ __device__ double1 _pixMakeZero<double1>() {return make_double1(0.);}
129 template<> inline __host__ __device__ double3 _pixMakeZero<double3>() {return make_double3(0.,0.,0.);}
130 template<> inline __host__ __device__ double4 _pixMakeZero<double4>() {return make_double4(0.,0.,0.,0.);}
131 
132 static inline __host__ __device__ uchar1 _pixMake(Ncv8u x) {return make_uchar1(x);}
133 static inline __host__ __device__ uchar3 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z) {return make_uchar3(x,y,z);}
134 static inline __host__ __device__ uchar4 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z, Ncv8u w) {return make_uchar4(x,y,z,w);}
135 static inline __host__ __device__ ushort1 _pixMake(Ncv16u x) {return make_ushort1(x);}
136 static inline __host__ __device__ ushort3 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z) {return make_ushort3(x,y,z);}
137 static inline __host__ __device__ ushort4 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z, Ncv16u w) {return make_ushort4(x,y,z,w);}
138 static inline __host__ __device__ uint1 _pixMake(Ncv32u x) {return make_uint1(x);}
139 static inline __host__ __device__ uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);}
140 static inline __host__ __device__ uint4 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z, Ncv32u w) {return make_uint4(x,y,z,w);}
141 static inline __host__ __device__ float1 _pixMake(Ncv32f x) {return make_float1(x);}
142 static inline __host__ __device__ float3 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z) {return make_float3(x,y,z);}
143 static inline __host__ __device__ float4 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z, Ncv32f w) {return make_float4(x,y,z,w);}
144 static inline __host__ __device__ double1 _pixMake(Ncv64f x) {return make_double1(x);}
145 static inline __host__ __device__ double3 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z) {return make_double3(x,y,z);}
146 static inline __host__ __device__ double4 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z, Ncv64f w) {return make_double4(x,y,z,w);}
147 
148 
149 template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampZ_CN {static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix);};
150 
151 template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 1> {
152 static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
153 {
154  Tout out;
155  _TDemoteClampZ(pix.x, out.x);
156  return out;
157 }};
158 
159 template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 3> {
160 static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
161 {
162  Tout out;
163  _TDemoteClampZ(pix.x, out.x);
164  _TDemoteClampZ(pix.y, out.y);
165  _TDemoteClampZ(pix.z, out.z);
166  return out;
167 }};
168 
169 template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 4> {
170 static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
171 {
172  Tout out;
173  _TDemoteClampZ(pix.x, out.x);
174  _TDemoteClampZ(pix.y, out.y);
175  _TDemoteClampZ(pix.z, out.z);
176  _TDemoteClampZ(pix.w, out.w);
177  return out;
178 }};
179 
180 template<typename Tin, typename Tout> static inline __host__ __device__ Tout _pixDemoteClampZ(Tin &pix)
181 {
183 }
184 
185 
186 template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampNN_CN {static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix);};
187 
188 template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 1> {
189 static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
190 {
191  Tout out;
192  _TDemoteClampNN(pix.x, out.x);
193  return out;
194 }};
195 
196 template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 3> {
197 static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
198 {
199  Tout out;
200  _TDemoteClampNN(pix.x, out.x);
201  _TDemoteClampNN(pix.y, out.y);
202  _TDemoteClampNN(pix.z, out.z);
203  return out;
204 }};
205 
206 template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 4> {
207 static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
208 {
209  Tout out;
210  _TDemoteClampNN(pix.x, out.x);
211  _TDemoteClampNN(pix.y, out.y);
212  _TDemoteClampNN(pix.z, out.z);
213  _TDemoteClampNN(pix.w, out.w);
214  return out;
215 }};
216 
217 template<typename Tin, typename Tout> static inline __host__ __device__ Tout _pixDemoteClampNN(Tin &pix)
218 {
220 }
221 
222 
223 template<typename Tin, typename Tout, typename Tw, Ncv32u CN> struct __pixScale_CN {static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w);};
224 
225 template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 1> {
226 static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
227 {
228  Tout out;
229  using TBout = typename TConvVec2Base<Tout>::TBase;
230  out.x = (TBout)(pix.x * w);
231  return out;
232 }};
233 
234 template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 3> {
235 static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
236 {
237  Tout out;
238  using TBout = typename TConvVec2Base<Tout>::TBase;
239  out.x = (TBout)(pix.x * w);
240  out.y = (TBout)(pix.y * w);
241  out.z = (TBout)(pix.z * w);
242  return out;
243 }};
244 
245 template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 4> {
246 static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
247 {
248  Tout out;
249  using TBout = typename TConvVec2Base<Tout>::TBase;
250  out.x = (TBout)(pix.x * w);
251  out.y = (TBout)(pix.y * w);
252  out.z = (TBout)(pix.z * w);
253  out.w = (TBout)(pix.w * w);
254  return out;
255 }};
256 
257 template<typename Tin, typename Tout, typename Tw> static __host__ __device__ Tout _pixScale(Tin &pix, Tw w)
258 {
260 }
261 
262 
263 template<typename Tin, typename Tout, Ncv32u CN> struct __pixAdd_CN {static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2);};
264 
265 template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 1> {
266 static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
267 {
268  Tout out;
269  out.x = pix1.x + pix2.x;
270  return out;
271 }};
272 
273 template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 3> {
274 static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
275 {
276  Tout out;
277  out.x = pix1.x + pix2.x;
278  out.y = pix1.y + pix2.y;
279  out.z = pix1.z + pix2.z;
280  return out;
281 }};
282 
283 template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 4> {
284 static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
285 {
286  Tout out;
287  out.x = pix1.x + pix2.x;
288  out.y = pix1.y + pix2.y;
289  out.z = pix1.z + pix2.z;
290  out.w = pix1.w + pix2.w;
291  return out;
292 }};
293 
294 template<typename Tin, typename Tout> static __host__ __device__ Tout _pixAdd(Tout &pix1, Tin &pix2)
295 {
297 }
298 
299 
300 template<typename Tin, typename Tout, Ncv32u CN> struct __pixDist_CN {static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2);};
301 
302 template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 1> {
303 static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
304 {
305  return Tout(SQR(pix1.x - pix2.x));
306 }};
307 
308 template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 3> {
309 static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
310 {
311  return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z));
312 }};
313 
314 template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 4> {
315 static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
316 {
317  return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z) + SQR(pix1.w - pix2.w));
318 }};
319 
320 template<typename Tin, typename Tout> static __host__ __device__ Tout _pixDist(Tin &pix1, Tin &pix2)
321 {
323 }
324 
325 
326 template <typename T> struct TAccPixWeighted;
327 template<> struct TAccPixWeighted<uchar1> {using type = double1;};
328 template<> struct TAccPixWeighted<uchar3> {using type = double3;};
329 template<> struct TAccPixWeighted<uchar4> {using type = double4;};
330 template<> struct TAccPixWeighted<ushort1> {using type = double1;};
331 template<> struct TAccPixWeighted<ushort3> {using type = double3;};
332 template<> struct TAccPixWeighted<ushort4> {using type = double4;};
333 template<> struct TAccPixWeighted<float1> {using type = double1;};
334 template<> struct TAccPixWeighted<float3> {using type = double3;};
335 template<> struct TAccPixWeighted<float4> {using type = double4;};
336 
337 template<typename Tfrom> struct TAccPixDist {};
338 template<> struct TAccPixDist<uchar1> {using type = Ncv32u;};
339 template<> struct TAccPixDist<uchar3> {using type = Ncv32u;};
340 template<> struct TAccPixDist<uchar4> {using type = Ncv32u;};
341 template<> struct TAccPixDist<ushort1> {using type = Ncv32u;};
342 template<> struct TAccPixDist<ushort3> {using type = Ncv32u;};
343 template<> struct TAccPixDist<ushort4> {using type = Ncv32u;};
344 template<> struct TAccPixDist<float1> {using type = Ncv32f;};
345 template<> struct TAccPixDist<float3> {using type = Ncv32f;};
346 template<> struct TAccPixDist<float4> {using type = Ncv32f;};
347 
348 #endif //_ncv_pixel_operations_hpp_
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)