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(); 51 template<> inline __host__ __device__ Ncv8u _pixMaxVal<Ncv8u>() {return UCHAR_MAX;} 52 template<> inline __host__ __device__ Ncv16u _pixMaxVal<Ncv16u>() {return USHRT_MAX;} 53 template<> inline __host__ __device__ Ncv32u _pixMaxVal<Ncv32u>() {return UINT_MAX;} 54 template<> inline __host__ __device__ Ncv8s _pixMaxVal<Ncv8s>() {return SCHAR_MAX;} 55 template<> inline __host__ __device__ Ncv16s _pixMaxVal<Ncv16s>() {return SHRT_MAX;} 56 template<> inline __host__ __device__ Ncv32s _pixMaxVal<Ncv32s>() {return INT_MAX;} 57 template<> inline __host__ __device__ Ncv32f _pixMaxVal<Ncv32f>() {return FLT_MAX;} 58 template<> inline __host__ __device__ Ncv64f _pixMaxVal<Ncv64f>() {return DBL_MAX;} 59 60 template<typename TBase> inline __host__ __device__ TBase _pixMinVal(); 61 template<> inline __host__ __device__ Ncv8u _pixMinVal<Ncv8u>() {return 0;} 62 template<> inline __host__ __device__ Ncv16u _pixMinVal<Ncv16u>() {return 0;} 63 template<> inline __host__ __device__ Ncv32u _pixMinVal<Ncv32u>() {return 0;} 64 template<> inline __host__ __device__ Ncv8s _pixMinVal<Ncv8s>() {return SCHAR_MIN;} 65 template<> inline __host__ __device__ Ncv16s _pixMinVal<Ncv16s>() {return SHRT_MIN;} 66 template<> inline __host__ __device__ Ncv32s _pixMinVal<Ncv32s>() {return INT_MIN;} 67 template<> inline __host__ __device__ Ncv32f _pixMinVal<Ncv32f>() {return FLT_MIN;} 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 107 template<typename Tin> inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a);} 108 template<typename Tin> inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a, 0, USHRT_MAX);} 109 template<typename Tin> inline __host__ __device__ void _TDemoteClampZ(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a, 0, UINT_MAX);} 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 113 template<typename Tin> inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv8u &out) {out = (Ncv8u)CLAMP_0_255(a+0.5f);} 114 template<typename Tin> inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv16u &out) {out = (Ncv16u)CLAMP(a+0.5f, 0, USHRT_MAX);} 115 template<typename Tin> inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32u &out) {out = (Ncv32u)CLAMP(a+0.5f, 0, UINT_MAX);} 116 template<typename Tin> inline __host__ __device__ void _TDemoteClampNN(Tin &a, Ncv32f &out) {out = (Ncv32f)a;} 117 118 template<typename Tout> inline Tout _pixMakeZero(); 119 template<> inline __host__ __device__ uchar1 _pixMakeZero<uchar1>() {return make_uchar1(0);} 120 template<> inline __host__ __device__ uchar3 _pixMakeZero<uchar3>() {return make_uchar3(0,0,0);} 121 template<> inline __host__ __device__ uchar4 _pixMakeZero<uchar4>() {return make_uchar4(0,0,0,0);} 122 template<> inline __host__ __device__ ushort1 _pixMakeZero<ushort1>() {return make_ushort1(0);} 123 template<> inline __host__ __device__ ushort3 _pixMakeZero<ushort3>() {return make_ushort3(0,0,0);} 124 template<> inline __host__ __device__ ushort4 _pixMakeZero<ushort4>() {return make_ushort4(0,0,0,0);} 125 template<> inline __host__ __device__ uint1 _pixMakeZero<uint1>() {return make_uint1(0);} 126 template<> inline __host__ __device__ uint3 _pixMakeZero<uint3>() {return make_uint3(0,0,0);} 127 template<> inline __host__ __device__ uint4 _pixMakeZero<uint4>() {return make_uint4(0,0,0,0);} 128 template<> inline __host__ __device__ float1 _pixMakeZero<float1>() {return make_float1(0.f);} 129 template<> inline __host__ __device__ float3 _pixMakeZero<float3>() {return make_float3(0.f,0.f,0.f);} 130 template<> inline __host__ __device__ float4 _pixMakeZero<float4>() {return make_float4(0.f,0.f,0.f,0.f);} 131 template<> inline __host__ __device__ double1 _pixMakeZero<double1>() {return make_double1(0.);} 132 template<> inline __host__ __device__ double3 _pixMakeZero<double3>() {return make_double3(0.,0.,0.);} 133 template<> inline __host__ __device__ double4 _pixMakeZero<double4>() {return make_double4(0.,0.,0.,0.);} 134 135 static inline __host__ __device__ uchar1 _pixMake(Ncv8u x) {return make_uchar1(x);} 136 static inline __host__ __device__ uchar3 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z) {return make_uchar3(x,y,z);} 137 static inline __host__ __device__ uchar4 _pixMake(Ncv8u x, Ncv8u y, Ncv8u z, Ncv8u w) {return make_uchar4(x,y,z,w);} 138 static inline __host__ __device__ ushort1 _pixMake(Ncv16u x) {return make_ushort1(x);} 139 static inline __host__ __device__ ushort3 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z) {return make_ushort3(x,y,z);} 140 static inline __host__ __device__ ushort4 _pixMake(Ncv16u x, Ncv16u y, Ncv16u z, Ncv16u w) {return make_ushort4(x,y,z,w);} 141 static inline __host__ __device__ uint1 _pixMake(Ncv32u x) {return make_uint1(x);} 142 static inline __host__ __device__ uint3 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z) {return make_uint3(x,y,z);} 143 static inline __host__ __device__ uint4 _pixMake(Ncv32u x, Ncv32u y, Ncv32u z, Ncv32u w) {return make_uint4(x,y,z,w);} 144 static inline __host__ __device__ float1 _pixMake(Ncv32f x) {return make_float1(x);} 145 static inline __host__ __device__ float3 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z) {return make_float3(x,y,z);} 146 static inline __host__ __device__ float4 _pixMake(Ncv32f x, Ncv32f y, Ncv32f z, Ncv32f w) {return make_float4(x,y,z,w);} 147 static inline __host__ __device__ double1 _pixMake(Ncv64f x) {return make_double1(x);} 148 static inline __host__ __device__ double3 _pixMake(Ncv64f x, Ncv64f y, Ncv64f z) {return make_double3(x,y,z);} 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> { 155 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> { 163 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> { 173 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 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> { 192 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> { 200 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> { 210 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 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> { 229 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> { 238 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> { 249 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 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> { 269 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> { 277 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> { 287 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 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> { 306 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> { 312 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> { 318 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 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
