Point Cloud Library (PCL) 1.12.0
Loading...
Searching...
No Matches
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.h>
45#include <float.h>
46#include "NCV.hpp"
47
48template<typename TBase> inline __host__ __device__ TBase _pixMaxVal();
49template<> inline __host__ __device__ Ncv8u _pixMaxVal<Ncv8u>() {return UCHAR_MAX;}
50template<> inline __host__ __device__ Ncv16u _pixMaxVal<Ncv16u>() {return USHRT_MAX;}
51template<> inline __host__ __device__ Ncv32u _pixMaxVal<Ncv32u>() {return UINT_MAX;}
52template<> inline __host__ __device__ Ncv8s _pixMaxVal<Ncv8s>() {return SCHAR_MAX;}
53template<> inline __host__ __device__ Ncv16s _pixMaxVal<Ncv16s>() {return SHRT_MAX;}
54template<> inline __host__ __device__ Ncv32s _pixMaxVal<Ncv32s>() {return INT_MAX;}
55template<> inline __host__ __device__ Ncv32f _pixMaxVal<Ncv32f>() {return FLT_MAX;}
56template<> inline __host__ __device__ Ncv64f _pixMaxVal<Ncv64f>() {return DBL_MAX;}
57
58template<typename TBase> inline __host__ __device__ TBase _pixMinVal();
59template<> inline __host__ __device__ Ncv8u _pixMinVal<Ncv8u>() {return 0;}
60template<> inline __host__ __device__ Ncv16u _pixMinVal<Ncv16u>() {return 0;}
61template<> inline __host__ __device__ Ncv32u _pixMinVal<Ncv32u>() {return 0;}
62template<> inline __host__ __device__ Ncv8s _pixMinVal<Ncv8s>() {return SCHAR_MIN;}
63template<> inline __host__ __device__ Ncv16s _pixMinVal<Ncv16s>() {return SHRT_MIN;}
64template<> inline __host__ __device__ Ncv32s _pixMinVal<Ncv32s>() {return INT_MIN;}
65template<> inline __host__ __device__ Ncv32f _pixMinVal<Ncv32f>() {return FLT_MIN;}
66template<> inline __host__ __device__ Ncv64f _pixMinVal<Ncv64f>() {return DBL_MIN;}
67
68template<typename Tvec> struct TConvVec2Base;
69template<> struct TConvVec2Base<uchar1> {using TBase = Ncv8u;};
70template<> struct TConvVec2Base<uchar3> {using TBase = Ncv8u;};
71template<> struct TConvVec2Base<uchar4> {using TBase = Ncv8u;};
72template<> struct TConvVec2Base<ushort1> {using TBase = Ncv16u;};
73template<> struct TConvVec2Base<ushort3> {using TBase = Ncv16u;};
74template<> struct TConvVec2Base<ushort4> {using TBase = Ncv16u;};
75template<> struct TConvVec2Base<uint1> {using TBase = Ncv32u;};
76template<> struct TConvVec2Base<uint3> {using TBase = Ncv32u;};
77template<> struct TConvVec2Base<uint4> {using TBase = Ncv32u;};
78template<> struct TConvVec2Base<float1> {using TBase = Ncv32f;};
79template<> struct TConvVec2Base<float3> {using TBase = Ncv32f;};
80template<> struct TConvVec2Base<float4> {using TBase = Ncv32f;};
81template<> struct TConvVec2Base<double1> {using TBase = Ncv64f;};
82template<> struct TConvVec2Base<double3> {using TBase = Ncv64f;};
83template<> struct TConvVec2Base<double4> {using TBase = Ncv64f;};
84
85#define NC(T) (sizeof(T) / sizeof(TConvVec2Base<T>::TBase))
86
87template<typename TBase, Ncv32u NC> struct TConvBase2Vec;
88template<> struct TConvBase2Vec<Ncv8u, 1> {using TVec = uchar1;};
89template<> struct TConvBase2Vec<Ncv8u, 3> {using TVec = uchar3;};
90template<> struct TConvBase2Vec<Ncv8u, 4> {using TVec = uchar4;};
91template<> struct TConvBase2Vec<Ncv16u, 1> {using TVec = ushort1;};
92template<> struct TConvBase2Vec<Ncv16u, 3> {using TVec = ushort3;};
93template<> struct TConvBase2Vec<Ncv16u, 4> {using TVec = ushort4;};
94template<> struct TConvBase2Vec<Ncv32u, 1> {using TVec = uint1;};
95template<> struct TConvBase2Vec<Ncv32u, 3> {using TVec = uint3;};
96template<> struct TConvBase2Vec<Ncv32u, 4> {using TVec = uint4;};
97template<> struct TConvBase2Vec<Ncv32f, 1> {using TVec = float1;};
98template<> struct TConvBase2Vec<Ncv32f, 3> {using TVec = float3;};
99template<> struct TConvBase2Vec<Ncv32f, 4> {using TVec = float4;};
100template<> struct TConvBase2Vec<Ncv64f, 1> {using TVec = double1;};
101template<> struct TConvBase2Vec<Ncv64f, 3> {using TVec = double3;};
102template<> struct TConvBase2Vec<Ncv64f, 4> {using TVec = double4;};
103
104//TODO: consider using CUDA intrinsics to avoid branching
105template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a);};
106template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a, 0, USHRT_MAX);}
107template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a, 0, UINT_MAX);}
108template<typename Tin> static inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}
109
110//TODO: consider using CUDA intrinsics to avoid branching
111template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a+0.5f);}
112template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a+0.5f, 0, USHRT_MAX);}
113template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a+0.5f, 0, UINT_MAX);}
114template<typename Tin> static inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32f &out) {out = (Ncv32f)a;}
115
116template<typename Tout> inline Tout _pixMakeZero();
117template<> inline __host__ __device__ uchar1 _pixMakeZero<uchar1>() {return make_uchar1(0);}
118template<> inline __host__ __device__ uchar3 _pixMakeZero<uchar3>() {return make_uchar3(0,0,0);}
119template<> inline __host__ __device__ uchar4 _pixMakeZero<uchar4>() {return make_uchar4(0,0,0,0);}
120template<> inline __host__ __device__ ushort1 _pixMakeZero<ushort1>() {return make_ushort1(0);}
121template<> inline __host__ __device__ ushort3 _pixMakeZero<ushort3>() {return make_ushort3(0,0,0);}
122template<> inline __host__ __device__ ushort4 _pixMakeZero<ushort4>() {return make_ushort4(0,0,0,0);}
123template<> inline __host__ __device__ uint1 _pixMakeZero<uint1>() {return make_uint1(0);}
124template<> inline __host__ __device__ uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,0);}
125template<> inline __host__ __device__ uint4 _pixMakeZero<uint4>() {return make_uint4(0,0,0,0);}
126template<> inline __host__ __device__ float1 _pixMakeZero<float1>() {return make_float1(0.f);}
127template<> inline __host__ __device__ float3 _pixMakeZero<float3>() {return make_float3(0.f,0.f,0.f);}
128template<> inline __host__ __device__ float4 _pixMakeZero<float4>() {return make_float4(0.f,0.f,0.f,0.f);}
129template<> inline __host__ __device__ double1 _pixMakeZero<double1>() {return make_double1(0.);}
130template<> inline __host__ __device__ double3 _pixMakeZero<double3>() {return make_double3(0.,0.,0.);}
131template<> inline __host__ __device__ double4 _pixMakeZero<double4>() {return make_double4(0.,0.,0.,0.);}
132
133static inline __host__ __device__ uchar1 _pixMake(Ncv8u x) {return make_uchar1(x);}
134static inline __host__ __device__ uchar3 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z) {return make_uchar3(x,y,z);}
135static inline __host__ __device__ uchar4 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z, Ncv8u w) {return make_uchar4(x,y,z,w);}
136static inline __host__ __device__ ushort1 _pixMake(Ncv16u x) {return make_ushort1(x);}
137static inline __host__ __device__ ushort3 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z) {return make_ushort3(x,y,z);}
138static inline __host__ __device__ ushort4 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z, Ncv16u w) {return make_ushort4(x,y,z,w);}
139static inline __host__ __device__ uint1 _pixMake(Ncv32u x) {return make_uint1(x);}
140static inline __host__ __device__ uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);}
141static inline __host__ __device__ uint4 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z, Ncv32u w) {return make_uint4(x,y,z,w);}
142static inline __host__ __device__ float1 _pixMake(Ncv32f x) {return make_float1(x);}
143static inline __host__ __device__ float3 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z) {return make_float3(x,y,z);}
144static inline __host__ __device__ float4 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z, Ncv32f w) {return make_float4(x,y,z,w);}
145static inline __host__ __device__ double1 _pixMake(Ncv64f x) {return make_double1(x);}
146static inline __host__ __device__ double3 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z) {return make_double3(x,y,z);}
147static inline __host__ __device__ double4 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z, Ncv64f w) {return make_double4(x,y,z,w);}
148
149
150template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampZ_CN {static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix);};
151
152template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 1> {
153static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
154{
155 Tout out;
156 _TDemoteClampZ(pix.x, out.x);
157 return out;
158}};
159
160template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 3> {
161static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
162{
163 Tout out;
164 _TDemoteClampZ(pix.x, out.x);
165 _TDemoteClampZ(pix.y, out.y);
166 _TDemoteClampZ(pix.z, out.z);
167 return out;
168}};
169
170template<typename Tin, typename Tout> struct __pixDemoteClampZ_CN<Tin, Tout, 4> {
171static __host__ __device__ Tout _pixDemoteClampZ_CN(Tin &pix)
172{
173 Tout out;
174 _TDemoteClampZ(pix.x, out.x);
175 _TDemoteClampZ(pix.y, out.y);
176 _TDemoteClampZ(pix.z, out.z);
177 _TDemoteClampZ(pix.w, out.w);
178 return out;
179}};
180
181template<typename Tin, typename Tout> static inline __host__ __device__ Tout _pixDemoteClampZ(Tin &pix)
182{
183 return __pixDemoteClampZ_CN<Tin, Tout, NC(Tin)>::_pixDemoteClampZ_CN(pix);
184}
185
186
187template<typename Tin, typename Tout, Ncv32u CN> struct __pixDemoteClampNN_CN {static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix);};
188
189template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 1> {
190static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
191{
192 Tout out;
193 _TDemoteClampNN(pix.x, out.x);
194 return out;
195}};
196
197template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 3> {
198static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
199{
200 Tout out;
201 _TDemoteClampNN(pix.x, out.x);
202 _TDemoteClampNN(pix.y, out.y);
203 _TDemoteClampNN(pix.z, out.z);
204 return out;
205}};
206
207template<typename Tin, typename Tout> struct __pixDemoteClampNN_CN<Tin, Tout, 4> {
208static __host__ __device__ Tout _pixDemoteClampNN_CN(Tin &pix)
209{
210 Tout out;
211 _TDemoteClampNN(pix.x, out.x);
212 _TDemoteClampNN(pix.y, out.y);
213 _TDemoteClampNN(pix.z, out.z);
214 _TDemoteClampNN(pix.w, out.w);
215 return out;
216}};
217
218template<typename Tin, typename Tout> static inline __host__ __device__ Tout _pixDemoteClampNN(Tin &pix)
219{
220 return __pixDemoteClampNN_CN<Tin, Tout, NC(Tin)>::_pixDemoteClampNN_CN(pix);
221}
222
223
224template<typename Tin, typename Tout, typename Tw, Ncv32u CN> struct __pixScale_CN {static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w);};
225
226template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 1> {
227static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
228{
229 Tout out;
230 using TBout = typename TConvVec2Base<Tout>::TBase;
231 out.x = (TBout)(pix.x * w);
232 return out;
233}};
234
235template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 3> {
236static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
237{
238 Tout out;
239 using TBout = typename TConvVec2Base<Tout>::TBase;
240 out.x = (TBout)(pix.x * w);
241 out.y = (TBout)(pix.y * w);
242 out.z = (TBout)(pix.z * w);
243 return out;
244}};
245
246template<typename Tin, typename Tout, typename Tw> struct __pixScale_CN<Tin, Tout, Tw, 4> {
247static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
248{
249 Tout out;
250 using TBout = typename TConvVec2Base<Tout>::TBase;
251 out.x = (TBout)(pix.x * w);
252 out.y = (TBout)(pix.y * w);
253 out.z = (TBout)(pix.z * w);
254 out.w = (TBout)(pix.w * w);
255 return out;
256}};
257
258template<typename Tin, typename Tout, typename Tw> static __host__ __device__ Tout _pixScale(Tin &pix, Tw w)
259{
260 return __pixScale_CN<Tin, Tout, Tw, NC(Tin)>::_pixScale_CN(pix, w);
261}
262
263
264template<typename Tin, typename Tout, Ncv32u CN> struct __pixAdd_CN {static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2);};
265
266template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 1> {
267static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
268{
269 Tout out;
270 out.x = pix1.x + pix2.x;
271 return out;
272}};
273
274template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 3> {
275static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
276{
277 Tout out;
278 out.x = pix1.x + pix2.x;
279 out.y = pix1.y + pix2.y;
280 out.z = pix1.z + pix2.z;
281 return out;
282}};
283
284template<typename Tin, typename Tout> struct __pixAdd_CN<Tin, Tout, 4> {
285static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
286{
287 Tout out;
288 out.x = pix1.x + pix2.x;
289 out.y = pix1.y + pix2.y;
290 out.z = pix1.z + pix2.z;
291 out.w = pix1.w + pix2.w;
292 return out;
293}};
294
295template<typename Tin, typename Tout> static __host__ __device__ Tout _pixAdd(Tout &pix1, Tin &pix2)
296{
297 return __pixAdd_CN<Tin, Tout, NC(Tin)>::_pixAdd_CN(pix1, pix2);
298}
299
300
301template<typename Tin, typename Tout, Ncv32u CN> struct __pixDist_CN {static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2);};
302
303template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 1> {
304static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
305{
306 return Tout(SQR(pix1.x - pix2.x));
307}};
308
309template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 3> {
310static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
311{
312 return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z));
313}};
314
315template<typename Tin, typename Tout> struct __pixDist_CN<Tin, Tout, 4> {
316static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
317{
318 return Tout(SQR(pix1.x - pix2.x) + SQR(pix1.y - pix2.y) + SQR(pix1.z - pix2.z) + SQR(pix1.w - pix2.w));
319}};
320
321template<typename Tin, typename Tout> static __host__ __device__ Tout _pixDist(Tin &pix1, Tin &pix2)
322{
323 return __pixDist_CN<Tin, Tout, NC(Tin)>::_pixDist_CN(pix1, pix2);
324}
325
326
327template <typename T> struct TAccPixWeighted;
328template<> struct TAccPixWeighted<uchar1> {using type = double1;};
329template<> struct TAccPixWeighted<uchar3> {using type = double3;};
330template<> struct TAccPixWeighted<uchar4> {using type = double4;};
331template<> struct TAccPixWeighted<ushort1> {using type = double1;};
332template<> struct TAccPixWeighted<ushort3> {using type = double3;};
333template<> struct TAccPixWeighted<ushort4> {using type = double4;};
334template<> struct TAccPixWeighted<float1> {using type = double1;};
335template<> struct TAccPixWeighted<float3> {using type = double3;};
336template<> struct TAccPixWeighted<float4> {using type = double4;};
337
338template<typename Tfrom> struct TAccPixDist {};
339template<> struct TAccPixDist<uchar1> {using type = Ncv32u;};
340template<> struct TAccPixDist<uchar3> {using type = Ncv32u;};
341template<> struct TAccPixDist<uchar4> {using type = Ncv32u;};
342template<> struct TAccPixDist<ushort1> {using type = Ncv32u;};
343template<> struct TAccPixDist<ushort3> {using type = Ncv32u;};
344template<> struct TAccPixDist<ushort4> {using type = Ncv32u;};
345template<> struct TAccPixDist<float1> {using type = Ncv32f;};
346template<> struct TAccPixDist<float3> {using type = Ncv32f;};
347template<> struct TAccPixDist<float4> {using type = Ncv32f;};
348
349#endif //_ncv_pixel_operations_hpp_
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
static __host__ __device__ Tout _pixAdd_CN(Tout &pix1, Tin &pix2)
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 _pixDemoteClampNN_CN(Tin &pix)
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 _pixDemoteClampZ_CN(Tin &pix)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixDist_CN(Tin &pix1, Tin &pix2)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)
static __host__ __device__ Tout _pixScale_CN(Tin &pix, Tw w)