vec_traits.hpp

00001 /*M///////////////////////////////////////////////////////////////////////////////////////
00002 //
00003 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
00004 //
00005 //  By downloading, copying, installing or using the software you agree to this license.
00006 //  If you do not agree to this license, do not download, install,
00007 //  copy or use the software.
00008 //
00009 //
00010 //                           License Agreement
00011 //                For Open Source Computer Vision Library
00012 //
00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
00015 // Third party copyrights are property of their respective owners.
00016 //
00017 // Redistribution and use in source and binary forms, with or without modification,
00018 // are permitted provided that the following conditions are met:
00019 //
00020 //   * Redistribution's of source code must retain the above copyright notice,
00021 //     this list of conditions and the following disclaimer.
00022 //
00023 //   * Redistribution's in binary form must reproduce the above copyright notice,
00024 //     this list of conditions and the following disclaimer in the documentation
00025 //     and/or other materials provided with the distribution.
00026 //
00027 //   * The name of the copyright holders may not be used to endorse or promote products
00028 //     derived from this software without specific prior written permission.
00029 //
00030 // This software is provided by the copyright holders and contributors "as is" and
00031 // any express or implied warranties, including, but not limited to, the implied
00032 // warranties of merchantability and fitness for a particular purpose are disclaimed.
00033 // In no event shall the Intel Corporation or contributors be liable for any direct,
00034 // indirect, incidental, special, exemplary, or consequential damages
00035 // (including, but not limited to, procurement of substitute goods or services;
00036 // loss of use, data, or profits; or business interruption) however caused
00037 // and on any theory of liability, whether in contract, strict liability,
00038 // or tort (including negligence or otherwise) arising in any way out of
00039 // the use of this software, even if advised of the possibility of such damage.
00040 //
00041 //M*/
00042 
00043 #ifndef __OPENCV_CUDA_VEC_TRAITS_HPP__
00044 #define __OPENCV_CUDA_VEC_TRAITS_HPP__
00045 
00046 #include "common.hpp "
00047 
00048 /** @file
00049  * @deprecated Use @ref cudev instead.
00050  */
00051 
00052 //! @cond IGNORED
00053 
00054 namespace cv { namespace cuda { namespace device
00055 {
00056     template<typename T, int N> struct TypeVec;
00057 
00058     struct __align__(8) uchar8
00059     {
00060         uchar a0, a1, a2, a3, a4, a5, a6, a7;
00061     };
00062     static __host__ __device__ __forceinline__ uchar8 make_uchar8(uchar a0, uchar a1, uchar a2, uchar a3, uchar a4, uchar a5, uchar a6, uchar a7)
00063     {
00064         uchar8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00065         return val;
00066     }
00067     struct __align__(8) char8
00068     {
00069         schar a0, a1, a2, a3, a4, a5, a6, a7;
00070     };
00071     static __host__ __device__ __forceinline__ char8 make_char8(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7)
00072     {
00073         char8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00074         return val;
00075     }
00076     struct __align__(16) ushort8
00077     {
00078         ushort a0, a1, a2, a3, a4, a5, a6, a7;
00079     };
00080     static __host__ __device__ __forceinline__ ushort8 make_ushort8(ushort a0, ushort a1, ushort a2, ushort a3, ushort a4, ushort a5, ushort a6, ushort a7)
00081     {
00082         ushort8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00083         return val;
00084     }
00085     struct __align__(16) short8
00086     {
00087         short a0, a1, a2, a3, a4, a5, a6, a7;
00088     };
00089     static __host__ __device__ __forceinline__ short8 make_short8(short a0, short a1, short a2, short a3, short a4, short a5, short a6, short a7)
00090     {
00091         short8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00092         return val;
00093     }
00094     struct __align__(32) uint8
00095     {
00096         uint a0, a1, a2, a3, a4, a5, a6, a7;
00097     };
00098     static __host__ __device__ __forceinline__ uint8 make_uint8(uint a0, uint a1, uint a2, uint a3, uint a4, uint a5, uint a6, uint a7)
00099     {
00100         uint8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00101         return val;
00102     }
00103     struct __align__(32) int8
00104     {
00105         int a0, a1, a2, a3, a4, a5, a6, a7;
00106     };
00107     static __host__ __device__ __forceinline__ int8 make_int8(int a0, int a1, int a2, int a3, int a4, int a5, int a6, int a7)
00108     {
00109         int8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00110         return val;
00111     }
00112     struct __align__(32) float8
00113     {
00114         float a0, a1, a2, a3, a4, a5, a6, a7;
00115     };
00116     static __host__ __device__ __forceinline__ float8 make_float8(float a0, float a1, float a2, float a3, float a4, float a5, float a6, float a7)
00117     {
00118         float8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00119         return val;
00120     }
00121     struct double8
00122     {
00123         double a0, a1, a2, a3, a4, a5, a6, a7;
00124     };
00125     static __host__ __device__ __forceinline__ double8 make_double8(double a0, double a1, double a2, double a3, double a4, double a5, double a6, double a7)
00126     {
00127         double8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
00128         return val;
00129     }
00130 
00131 #define OPENCV_CUDA_IMPLEMENT_TYPE_VEC(type) \
00132     template<> struct TypeVec<type, 1> { typedef type vec_type; }; \
00133     template<> struct TypeVec<type ## 1, 1> { typedef type ## 1 vec_type; }; \
00134     template<> struct TypeVec<type, 2> { typedef type ## 2 vec_type; }; \
00135     template<> struct TypeVec<type ## 2, 2> { typedef type ## 2 vec_type; }; \
00136     template<> struct TypeVec<type, 3> { typedef type ## 3 vec_type; }; \
00137     template<> struct TypeVec<type ## 3, 3> { typedef type ## 3 vec_type; }; \
00138     template<> struct TypeVec<type, 4> { typedef type ## 4 vec_type; }; \
00139     template<> struct TypeVec<type ## 4, 4> { typedef type ## 4 vec_type; }; \
00140     template<> struct TypeVec<type, 8> { typedef type ## 8 vec_type; }; \
00141     template<> struct TypeVec<type ## 8, 8> { typedef type ## 8 vec_type; };
00142 
00143     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uchar)
00144     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(char)
00145     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(ushort)
00146     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(short)
00147     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(int)
00148     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uint)
00149     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(float)
00150     OPENCV_CUDA_IMPLEMENT_TYPE_VEC(double)
00151 
00152     #undef OPENCV_CUDA_IMPLEMENT_TYPE_VEC
00153 
00154     template<> struct TypeVec<schar, 1> { typedef schar vec_type; };
00155     template<> struct TypeVec<schar, 2> { typedef char2 vec_type; };
00156     template<> struct TypeVec<schar, 3> { typedef char3 vec_type; };
00157     template<> struct TypeVec<schar, 4> { typedef char4 vec_type; };
00158     template<> struct TypeVec<schar, 8> { typedef char8 vec_type; };
00159 
00160     template<> struct TypeVec<bool, 1> { typedef uchar vec_type; };
00161     template<> struct TypeVec<bool, 2> { typedef uchar2 vec_type; };
00162     template<> struct TypeVec<bool, 3> { typedef uchar3 vec_type; };
00163     template<> struct TypeVec<bool, 4> { typedef uchar4 vec_type; };
00164     template<> struct TypeVec<bool, 8> { typedef uchar8 vec_type; };
00165 
00166     template<typename T> struct VecTraits;
00167 
00168 #define OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(type) \
00169     template<> struct VecTraits<type> \
00170     { \
00171         typedef type elem_type; \
00172         enum {cn=1}; \
00173         static __device__ __host__ __forceinline__ type all(type v) {return v;} \
00174         static __device__ __host__ __forceinline__ type make(type x) {return x;} \
00175         static __device__ __host__ __forceinline__ type make(const type* v) {return *v;} \
00176     }; \
00177     template<> struct VecTraits<type ## 1> \
00178     { \
00179         typedef type elem_type; \
00180         enum {cn=1}; \
00181         static __device__ __host__ __forceinline__ type ## 1 all(type v) {return make_ ## type ## 1(v);} \
00182         static __device__ __host__ __forceinline__ type ## 1 make(type x) {return make_ ## type ## 1(x);} \
00183         static __device__ __host__ __forceinline__ type ## 1 make(const type* v) {return make_ ## type ## 1(*v);} \
00184     }; \
00185     template<> struct VecTraits<type ## 2> \
00186     { \
00187         typedef type elem_type; \
00188         enum {cn=2}; \
00189         static __device__ __host__ __forceinline__ type ## 2 all(type v) {return make_ ## type ## 2(v, v);} \
00190         static __device__ __host__ __forceinline__ type ## 2 make(type x, type y) {return make_ ## type ## 2(x, y);} \
00191         static __device__ __host__ __forceinline__ type ## 2 make(const type* v) {return make_ ## type ## 2(v[0], v[1]);} \
00192     }; \
00193     template<> struct VecTraits<type ## 3> \
00194     { \
00195         typedef type elem_type; \
00196         enum {cn=3}; \
00197         static __device__ __host__ __forceinline__ type ## 3 all(type v) {return make_ ## type ## 3(v, v, v);} \
00198         static __device__ __host__ __forceinline__ type ## 3 make(type x, type y, type z) {return make_ ## type ## 3(x, y, z);} \
00199         static __device__ __host__ __forceinline__ type ## 3 make(const type* v) {return make_ ## type ## 3(v[0], v[1], v[2]);} \
00200     }; \
00201     template<> struct VecTraits<type ## 4> \
00202     { \
00203         typedef type elem_type; \
00204         enum {cn=4}; \
00205         static __device__ __host__ __forceinline__ type ## 4 all(type v) {return make_ ## type ## 4(v, v, v, v);} \
00206         static __device__ __host__ __forceinline__ type ## 4 make(type x, type y, type z, type w) {return make_ ## type ## 4(x, y, z, w);} \
00207         static __device__ __host__ __forceinline__ type ## 4 make(const type* v) {return make_ ## type ## 4(v[0], v[1], v[2], v[3]);} \
00208     }; \
00209     template<> struct VecTraits<type ## 8> \
00210     { \
00211         typedef type elem_type; \
00212         enum {cn=8}; \
00213         static __device__ __host__ __forceinline__ type ## 8 all(type v) {return make_ ## type ## 8(v, v, v, v, v, v, v, v);} \
00214         static __device__ __host__ __forceinline__ type ## 8 make(type a0, type a1, type a2, type a3, type a4, type a5, type a6, type a7) {return make_ ## type ## 8(a0, a1, a2, a3, a4, a5, a6, a7);} \
00215         static __device__ __host__ __forceinline__ type ## 8 make(const type* v) {return make_ ## type ## 8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);} \
00216     };
00217 
00218     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uchar)
00219     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(ushort)
00220     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(short)
00221     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(int)
00222     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uint)
00223     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(float)
00224     OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(double)
00225 
00226     #undef OPENCV_CUDA_IMPLEMENT_VEC_TRAITS
00227 
00228     template<> struct VecTraits<char>
00229     {
00230         typedef char elem_type;
00231         enum {cn=1};
00232         static __device__ __host__ __forceinline__ char all(char v) {return v;}
00233         static __device__ __host__ __forceinline__ char make(char x) {return x;}
00234         static __device__ __host__ __forceinline__ char make(const char* x) {return *x;}
00235     };
00236     template<> struct VecTraits<schar>
00237     {
00238         typedef schar elem_type;
00239         enum {cn=1};
00240         static __device__ __host__ __forceinline__ schar all(schar v) {return v;}
00241         static __device__ __host__ __forceinline__ schar make(schar x) {return x;}
00242         static __device__ __host__ __forceinline__ schar make(const schar* x) {return *x;}
00243     };
00244     template<> struct VecTraits<char1>
00245     {
00246         typedef schar elem_type;
00247         enum {cn=1};
00248         static __device__ __host__ __forceinline__ char1 all(schar v) {return make_char1(v);}
00249         static __device__ __host__ __forceinline__ char1 make(schar x) {return make_char1(x);}
00250         static __device__ __host__ __forceinline__ char1 make(const schar* v) {return make_char1(v[0]);}
00251     };
00252     template<> struct VecTraits<char2>
00253     {
00254         typedef schar elem_type;
00255         enum {cn=2};
00256         static __device__ __host__ __forceinline__ char2 all(schar v) {return make_char2(v, v);}
00257         static __device__ __host__ __forceinline__ char2 make(schar x, schar y) {return make_char2(x, y);}
00258         static __device__ __host__ __forceinline__ char2 make(const schar* v) {return make_char2(v[0], v[1]);}
00259     };
00260     template<> struct VecTraits<char3>
00261     {
00262         typedef schar elem_type;
00263         enum {cn=3};
00264         static __device__ __host__ __forceinline__ char3 all(schar v) {return make_char3(v, v, v);}
00265         static __device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z) {return make_char3(x, y, z);}
00266         static __device__ __host__ __forceinline__ char3 make(const schar* v) {return make_char3(v[0], v[1], v[2]);}
00267     };
00268     template<> struct VecTraits<char4>
00269     {
00270         typedef schar elem_type;
00271         enum {cn=4};
00272         static __device__ __host__ __forceinline__ char4 all(schar v) {return make_char4(v, v, v, v);}
00273         static __device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w) {return make_char4(x, y, z, w);}
00274         static __device__ __host__ __forceinline__ char4 make(const schar* v) {return make_char4(v[0], v[1], v[2], v[3]);}
00275     };
00276     template<> struct VecTraits<char8>
00277     {
00278         typedef schar elem_type;
00279         enum {cn=8};
00280         static __device__ __host__ __forceinline__ char8 all(schar v) {return make_char8(v, v, v, v, v, v, v, v);}
00281         static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) {return make_char8(a0, a1, a2, a3, a4, a5, a6, a7);}
00282         static __device__ __host__ __forceinline__ char8 make(const schar* v) {return make_char8(v[0], v[1], v[2], v[3], v[4], v[5], v[6], v[7]);}
00283     };
00284 }}} // namespace cv { namespace cuda { namespace cudev
00285 
00286 //! @endcond
00287 
00288 #endif // __OPENCV_CUDA_VEC_TRAITS_HPP__
00289