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 OPENCV_CUDA_VEC_TRAITS_HPP 44 #define OPENCV_CUDA_VEC_TRAITS_HPP 45 46 #include "common.hpp" 47 48 /** @file 49 * @deprecated Use @ref cudev instead. 50 */ 51 52 //! @cond IGNORED 53 54 namespace cv { namespace cuda { namespace device 55 { 56 template<typename T, int N> struct TypeVec; 57 58 struct __align__(8) uchar8 59 { 60 uchar a0, a1, a2, a3, a4, a5, a6, a7; 61 }; make_uchar8(uchar a0,uchar a1,uchar a2,uchar a3,uchar a4,uchar a5,uchar a6,uchar a7)62 static __host__ __device__ __forceinline__ uchar8 make_uchar8(uchar a0, uchar a1, uchar a2, uchar a3, uchar a4, uchar a5, uchar a6, uchar a7) 63 { 64 uchar8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 65 return val; 66 } 67 struct __align__(8) char8 68 { 69 schar a0, a1, a2, a3, a4, a5, a6, a7; 70 }; make_char8(schar a0,schar a1,schar a2,schar a3,schar a4,schar a5,schar a6,schar a7)71 static __host__ __device__ __forceinline__ char8 make_char8(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7) 72 { 73 char8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 74 return val; 75 } 76 struct __align__(16) ushort8 77 { 78 ushort a0, a1, a2, a3, a4, a5, a6, a7; 79 }; make_ushort8(ushort a0,ushort a1,ushort a2,ushort a3,ushort a4,ushort a5,ushort a6,ushort a7)80 static __host__ __device__ __forceinline__ ushort8 make_ushort8(ushort a0, ushort a1, ushort a2, ushort a3, ushort a4, ushort a5, ushort a6, ushort a7) 81 { 82 ushort8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 83 return val; 84 } 85 struct __align__(16) short8 86 { 87 short a0, a1, a2, a3, a4, a5, a6, a7; 88 }; make_short8(short a0,short a1,short a2,short a3,short a4,short a5,short a6,short a7)89 static __host__ __device__ __forceinline__ short8 make_short8(short a0, short a1, short a2, short a3, short a4, short a5, short a6, short a7) 90 { 91 short8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 92 return val; 93 } 94 struct __align__(32) uint8 95 { 96 uint a0, a1, a2, a3, a4, a5, a6, a7; 97 }; make_uint8(uint a0,uint a1,uint a2,uint a3,uint a4,uint a5,uint a6,uint a7)98 static __host__ __device__ __forceinline__ uint8 make_uint8(uint a0, uint a1, uint a2, uint a3, uint a4, uint a5, uint a6, uint a7) 99 { 100 uint8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 101 return val; 102 } 103 struct __align__(32) int8 104 { 105 int a0, a1, a2, a3, a4, a5, a6, a7; 106 }; make_int8(int a0,int a1,int a2,int a3,int a4,int a5,int a6,int a7)107 static __host__ __device__ __forceinline__ int8 make_int8(int a0, int a1, int a2, int a3, int a4, int a5, int a6, int a7) 108 { 109 int8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 110 return val; 111 } 112 struct __align__(32) float8 113 { 114 float a0, a1, a2, a3, a4, a5, a6, a7; 115 }; make_float8(float a0,float a1,float a2,float a3,float a4,float a5,float a6,float a7)116 static __host__ __device__ __forceinline__ float8 make_float8(float a0, float a1, float a2, float a3, float a4, float a5, float a6, float a7) 117 { 118 float8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 119 return val; 120 } 121 struct double8 122 { 123 double a0, a1, a2, a3, a4, a5, a6, a7; 124 }; make_double8(double a0,double a1,double a2,double a3,double a4,double a5,double a6,double a7)125 static __host__ __device__ __forceinline__ double8 make_double8(double a0, double a1, double a2, double a3, double a4, double a5, double a6, double a7) 126 { 127 double8 val = {a0, a1, a2, a3, a4, a5, a6, a7}; 128 return val; 129 } 130 131 #define OPENCV_CUDA_IMPLEMENT_TYPE_VEC(type) \ 132 template<> struct TypeVec<type, 1> { typedef type vec_type; }; \ 133 template<> struct TypeVec<type ## 1, 1> { typedef type ## 1 vec_type; }; \ 134 template<> struct TypeVec<type, 2> { typedef type ## 2 vec_type; }; \ 135 template<> struct TypeVec<type ## 2, 2> { typedef type ## 2 vec_type; }; \ 136 template<> struct TypeVec<type, 3> { typedef type ## 3 vec_type; }; \ 137 template<> struct TypeVec<type ## 3, 3> { typedef type ## 3 vec_type; }; \ 138 template<> struct TypeVec<type, 4> { typedef type ## 4 vec_type; }; \ 139 template<> struct TypeVec<type ## 4, 4> { typedef type ## 4 vec_type; }; \ 140 template<> struct TypeVec<type, 8> { typedef type ## 8 vec_type; }; \ 141 template<> struct TypeVec<type ## 8, 8> { typedef type ## 8 vec_type; }; 142 143 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uchar) 144 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(char) 145 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(ushort) 146 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(short) 147 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(int) 148 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(uint) 149 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(float) 150 OPENCV_CUDA_IMPLEMENT_TYPE_VEC(double) 151 152 #undef OPENCV_CUDA_IMPLEMENT_TYPE_VEC 153 154 template<> struct TypeVec<schar, 1> { typedef schar vec_type; }; 155 template<> struct TypeVec<schar, 2> { typedef char2 vec_type; }; 156 template<> struct TypeVec<schar, 3> { typedef char3 vec_type; }; 157 template<> struct TypeVec<schar, 4> { typedef char4 vec_type; }; 158 template<> struct TypeVec<schar, 8> { typedef char8 vec_type; }; 159 160 template<> struct TypeVec<bool, 1> { typedef uchar vec_type; }; 161 template<> struct TypeVec<bool, 2> { typedef uchar2 vec_type; }; 162 template<> struct TypeVec<bool, 3> { typedef uchar3 vec_type; }; 163 template<> struct TypeVec<bool, 4> { typedef uchar4 vec_type; }; 164 template<> struct TypeVec<bool, 8> { typedef uchar8 vec_type; }; 165 166 template<typename T> struct VecTraits; 167 168 #define OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(type) \ 169 template<> struct VecTraits<type> \ 170 { \ 171 typedef type elem_type; \ 172 enum {cn=1}; \ 173 static __device__ __host__ __forceinline__ type all(type v) {return v;} \ 174 static __device__ __host__ __forceinline__ type make(type x) {return x;} \ 175 static __device__ __host__ __forceinline__ type make(const type* v) {return *v;} \ 176 }; \ 177 template<> struct VecTraits<type ## 1> \ 178 { \ 179 typedef type elem_type; \ 180 enum {cn=1}; \ 181 static __device__ __host__ __forceinline__ type ## 1 all(type v) {return make_ ## type ## 1(v);} \ 182 static __device__ __host__ __forceinline__ type ## 1 make(type x) {return make_ ## type ## 1(x);} \ 183 static __device__ __host__ __forceinline__ type ## 1 make(const type* v) {return make_ ## type ## 1(*v);} \ 184 }; \ 185 template<> struct VecTraits<type ## 2> \ 186 { \ 187 typedef type elem_type; \ 188 enum {cn=2}; \ 189 static __device__ __host__ __forceinline__ type ## 2 all(type v) {return make_ ## type ## 2(v, v);} \ 190 static __device__ __host__ __forceinline__ type ## 2 make(type x, type y) {return make_ ## type ## 2(x, y);} \ 191 static __device__ __host__ __forceinline__ type ## 2 make(const type* v) {return make_ ## type ## 2(v[0], v[1]);} \ 192 }; \ 193 template<> struct VecTraits<type ## 3> \ 194 { \ 195 typedef type elem_type; \ 196 enum {cn=3}; \ 197 static __device__ __host__ __forceinline__ type ## 3 all(type v) {return make_ ## type ## 3(v, v, v);} \ 198 static __device__ __host__ __forceinline__ type ## 3 make(type x, type y, type z) {return make_ ## type ## 3(x, y, z);} \ 199 static __device__ __host__ __forceinline__ type ## 3 make(const type* v) {return make_ ## type ## 3(v[0], v[1], v[2]);} \ 200 }; \ 201 template<> struct VecTraits<type ## 4> \ 202 { \ 203 typedef type elem_type; \ 204 enum {cn=4}; \ 205 static __device__ __host__ __forceinline__ type ## 4 all(type v) {return make_ ## type ## 4(v, v, v, v);} \ 206 static __device__ __host__ __forceinline__ type ## 4 make(type x, type y, type z, type w) {return make_ ## type ## 4(x, y, z, w);} \ 207 static __device__ __host__ __forceinline__ type ## 4 make(const type* v) {return make_ ## type ## 4(v[0], v[1], v[2], v[3]);} \ 208 }; \ 209 template<> struct VecTraits<type ## 8> \ 210 { \ 211 typedef type elem_type; \ 212 enum {cn=8}; \ 213 static __device__ __host__ __forceinline__ type ## 8 all(type v) {return make_ ## type ## 8(v, v, v, v, v, v, v, v);} \ 214 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);} \ 215 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]);} \ 216 }; 217 218 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uchar) 219 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(ushort) 220 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(short) 221 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(int) 222 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(uint) 223 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(float) 224 OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(double) 225 226 #undef OPENCV_CUDA_IMPLEMENT_VEC_TRAITS 227 228 template<> struct VecTraits<char> 229 { 230 typedef char elem_type; 231 enum {cn=1}; allcv::cuda::device::VecTraits232 static __device__ __host__ __forceinline__ char all(char v) {return v;} makecv::cuda::device::VecTraits233 static __device__ __host__ __forceinline__ char make(char x) {return x;} makecv::cuda::device::VecTraits234 static __device__ __host__ __forceinline__ char make(const char* x) {return *x;} 235 }; 236 template<> struct VecTraits<schar> 237 { 238 typedef schar elem_type; 239 enum {cn=1}; allcv::cuda::device::VecTraits240 static __device__ __host__ __forceinline__ schar all(schar v) {return v;} makecv::cuda::device::VecTraits241 static __device__ __host__ __forceinline__ schar make(schar x) {return x;} makecv::cuda::device::VecTraits242 static __device__ __host__ __forceinline__ schar make(const schar* x) {return *x;} 243 }; 244 template<> struct VecTraits<char1> 245 { 246 typedef schar elem_type; 247 enum {cn=1}; allcv::cuda::device::VecTraits248 static __device__ __host__ __forceinline__ char1 all(schar v) {return make_char1(v);} makecv::cuda::device::VecTraits249 static __device__ __host__ __forceinline__ char1 make(schar x) {return make_char1(x);} makecv::cuda::device::VecTraits250 static __device__ __host__ __forceinline__ char1 make(const schar* v) {return make_char1(v[0]);} 251 }; 252 template<> struct VecTraits<char2> 253 { 254 typedef schar elem_type; 255 enum {cn=2}; allcv::cuda::device::VecTraits256 static __device__ __host__ __forceinline__ char2 all(schar v) {return make_char2(v, v);} makecv::cuda::device::VecTraits257 static __device__ __host__ __forceinline__ char2 make(schar x, schar y) {return make_char2(x, y);} makecv::cuda::device::VecTraits258 static __device__ __host__ __forceinline__ char2 make(const schar* v) {return make_char2(v[0], v[1]);} 259 }; 260 template<> struct VecTraits<char3> 261 { 262 typedef schar elem_type; 263 enum {cn=3}; allcv::cuda::device::VecTraits264 static __device__ __host__ __forceinline__ char3 all(schar v) {return make_char3(v, v, v);} makecv::cuda::device::VecTraits265 static __device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z) {return make_char3(x, y, z);} makecv::cuda::device::VecTraits266 static __device__ __host__ __forceinline__ char3 make(const schar* v) {return make_char3(v[0], v[1], v[2]);} 267 }; 268 template<> struct VecTraits<char4> 269 { 270 typedef schar elem_type; 271 enum {cn=4}; allcv::cuda::device::VecTraits272 static __device__ __host__ __forceinline__ char4 all(schar v) {return make_char4(v, v, v, v);} makecv::cuda::device::VecTraits273 static __device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w) {return make_char4(x, y, z, w);} makecv::cuda::device::VecTraits274 static __device__ __host__ __forceinline__ char4 make(const schar* v) {return make_char4(v[0], v[1], v[2], v[3]);} 275 }; 276 template<> struct VecTraits<char8> 277 { 278 typedef schar elem_type; 279 enum {cn=8}; allcv::cuda::device::VecTraits280 static __device__ __host__ __forceinline__ char8 all(schar v) {return make_char8(v, v, v, v, v, v, v, v);} makecv::cuda::device::VecTraits281 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);} makecv::cuda::device::VecTraits282 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]);} 283 }; 284 }}} // namespace cv { namespace cuda { namespace cudev 285 286 //! @endcond 287 288 #endif // OPENCV_CUDA_VEC_TRAITS_HPP 289