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