1 /*M///////////////////////////////////////////////////////////////////////////////////////
\r
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
\r
5 // By downloading, copying, installing or using the software you agree to this license.
\r
6 // If you do not agree to this license, do not download, install,
\r
7 // copy or use the software.
\r
10 // License Agreement
\r
11 // For Open Source Computer Vision Library
\r
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
\r
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
\r
15 // Third party copyrights are property of their respective owners.
\r
17 // Redistribution and use in source and binary forms, with or without modification,
\r
18 // are permitted provided that the following conditions are met:
\r
20 // * Redistribution's of source code must retain the above copyright notice,
\r
21 // this list of conditions and the following disclaimer.
\r
23 // * Redistribution's in binary form must reproduce the above copyright notice,
\r
24 // this list of conditions and the following disclaimer in the documentation
\r
25 // and/or other materials provided with the distribution.
\r
27 // * The name of the copyright holders may not be used to endorse or promote products
\r
28 // derived from this software without specific prior written permission.
\r
30 // This software is provided by the copyright holders and contributors "as is" and
\r
31 // any express or implied warranties, including, but not limited to, the implied
\r
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
\r
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
\r
34 // indirect, incidental, special, exemplary, or consequential damages
\r
35 // (including, but not limited to, procurement of substitute goods or services;
\r
36 // loss of use, data, or profits; or business interruption) however caused
\r
37 // and on any theory of liability, whether in contract, strict liability,
\r
38 // or tort (including negligence or otherwise) arising in any way out of
\r
39 // the use of this software, even if advised of the possibility of such damage.
\r
43 #ifndef __OPENCV_GPU_VEC_TRAITS_HPP__
\r
44 #define __OPENCV_GPU_VEC_TRAITS_HPP__
\r
46 #include "internal_shared.hpp"
\r
48 namespace cv { namespace gpu { namespace device
\r
50 template<typename T, int N> struct TypeVec;
\r
52 struct __align__(8) uchar8
\r
54 uchar a0, a1, a2, a3, a4, a5, a6, a7;
\r
56 static __host__ __device__ __forceinline__ uchar8 make_uchar8(uchar a0, uchar a1, uchar a2, uchar a3, uchar a4, uchar a5, uchar a6, uchar a7)
\r
58 uchar8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
61 struct __align__(8) char8
\r
63 schar a0, a1, a2, a3, a4, a5, a6, a7;
\r
65 static __host__ __device__ __forceinline__ char8 make_char8(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7)
\r
67 char8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
70 struct __align__(16) ushort8
\r
72 ushort a0, a1, a2, a3, a4, a5, a6, a7;
\r
74 static __host__ __device__ __forceinline__ ushort8 make_ushort8(ushort a0, ushort a1, ushort a2, ushort a3, ushort a4, ushort a5, ushort a6, ushort a7)
\r
76 ushort8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
79 struct __align__(16) short8
\r
81 short a0, a1, a2, a3, a4, a5, a6, a7;
\r
83 static __host__ __device__ __forceinline__ short8 make_short8(short a0, short a1, short a2, short a3, short a4, short a5, short a6, short a7)
\r
85 short8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
88 struct __align__(32) uint8
\r
90 uint a0, a1, a2, a3, a4, a5, a6, a7;
\r
92 static __host__ __device__ __forceinline__ uint8 make_uint8(uint a0, uint a1, uint a2, uint a3, uint a4, uint a5, uint a6, uint a7)
\r
94 uint8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
97 struct __align__(32) int8
\r
99 int a0, a1, a2, a3, a4, a5, a6, a7;
\r
101 static __host__ __device__ __forceinline__ int8 make_int8(int a0, int a1, int a2, int a3, int a4, int a5, int a6, int a7)
\r
103 int8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
106 struct __align__(32) float8
\r
108 float a0, a1, a2, a3, a4, a5, a6, a7;
\r
110 static __host__ __device__ __forceinline__ float8 make_float8(float a0, float a1, float a2, float a3, float a4, float a5, float a6, float a7)
\r
112 float8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
117 double a0, a1, a2, a3, a4, a5, a6, a7;
\r
119 static __host__ __device__ __forceinline__ double8 make_double8(double a0, double a1, double a2, double a3, double a4, double a5, double a6, double a7)
\r
121 double8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
\r
125 #define OPENCV_GPU_IMPLEMENT_TYPE_VEC(type) \
\r
126 template<> struct TypeVec<type, 1> { typedef type vec_type; }; \
\r
127 template<> struct TypeVec<type ## 1, 1> { typedef type ## 1 vec_type; }; \
\r
128 template<> struct TypeVec<type, 2> { typedef type ## 2 vec_type; }; \
\r
129 template<> struct TypeVec<type ## 2, 2> { typedef type ## 2 vec_type; }; \
\r
130 template<> struct TypeVec<type, 3> { typedef type ## 3 vec_type; }; \
\r
131 template<> struct TypeVec<type ## 3, 3> { typedef type ## 3 vec_type; }; \
\r
132 template<> struct TypeVec<type, 4> { typedef type ## 4 vec_type; }; \
\r
133 template<> struct TypeVec<type ## 4, 4> { typedef type ## 4 vec_type; }; \
\r
134 template<> struct TypeVec<type, 8> { typedef type ## 8 vec_type; }; \
\r
135 template<> struct TypeVec<type ## 8, 8> { typedef type ## 8 vec_type; };
\r
137 OPENCV_GPU_IMPLEMENT_TYPE_VEC(uchar)
\r
138 OPENCV_GPU_IMPLEMENT_TYPE_VEC(char)
\r
139 OPENCV_GPU_IMPLEMENT_TYPE_VEC(ushort)
\r
140 OPENCV_GPU_IMPLEMENT_TYPE_VEC(short)
\r
141 OPENCV_GPU_IMPLEMENT_TYPE_VEC(int)
\r
142 OPENCV_GPU_IMPLEMENT_TYPE_VEC(uint)
\r
143 OPENCV_GPU_IMPLEMENT_TYPE_VEC(float)
\r
144 OPENCV_GPU_IMPLEMENT_TYPE_VEC(double)
\r
146 #undef OPENCV_GPU_IMPLEMENT_TYPE_VEC
\r
148 template<> struct TypeVec<schar, 1> { typedef schar vec_type; };
\r
149 template<> struct TypeVec<schar, 2> { typedef char2 vec_type; };
\r
150 template<> struct TypeVec<schar, 3> { typedef char3 vec_type; };
\r
151 template<> struct TypeVec<schar, 4> { typedef char4 vec_type; };
\r
152 template<> struct TypeVec<schar, 8> { typedef char8 vec_type; };
\r
154 template<> struct TypeVec<bool, 1> { typedef uchar vec_type; };
\r
155 template<> struct TypeVec<bool, 2> { typedef uchar2 vec_type; };
\r
156 template<> struct TypeVec<bool, 3> { typedef uchar3 vec_type; };
\r
157 template<> struct TypeVec<bool, 4> { typedef uchar4 vec_type; };
\r
158 template<> struct TypeVec<bool, 8> { typedef uchar8 vec_type; };
\r
160 template<typename T> struct VecTraits;
\r
162 #define OPENCV_GPU_IMPLEMENT_VEC_TRAITS(type) \
\r
163 template<> struct VecTraits<type> \
\r
165 typedef type elem_type; \
\r
167 static __device__ __host__ __forceinline__ type all(type v) {return v;} \
\r
168 static __device__ __host__ __forceinline__ type make(type x) {return x;} \
\r
169 static __device__ __host__ __forceinline__ type make(const type* v) {return *v;} \
\r
171 template<> struct VecTraits<type ## 1> \
\r
173 typedef type elem_type; \
\r
175 static __device__ __host__ __forceinline__ type ## 1 all(type v) {return make_ ## type ## 1(v);} \
\r
176 static __device__ __host__ __forceinline__ type ## 1 make(type x) {return make_ ## type ## 1(x);} \
\r
177 static __device__ __host__ __forceinline__ type ## 1 make(const type* v) {return make_ ## type ## 1(*v);} \
\r
179 template<> struct VecTraits<type ## 2> \
\r
181 typedef type elem_type; \
\r
183 static __device__ __host__ __forceinline__ type ## 2 all(type v) {return make_ ## type ## 2(v, v);} \
\r
184 static __device__ __host__ __forceinline__ type ## 2 make(type x, type y) {return make_ ## type ## 2(x, y);} \
\r
185 static __device__ __host__ __forceinline__ type ## 2 make(const type* v) {return make_ ## type ## 2(v[0], v[1]);} \
\r
187 template<> struct VecTraits<type ## 3> \
\r
189 typedef type elem_type; \
\r
191 static __device__ __host__ __forceinline__ type ## 3 all(type v) {return make_ ## type ## 3(v, v, v);} \
\r
192 static __device__ __host__ __forceinline__ type ## 3 make(type x, type y, type z) {return make_ ## type ## 3(x, y, z);} \
\r
193 static __device__ __host__ __forceinline__ type ## 3 make(const type* v) {return make_ ## type ## 3(v[0], v[1], v[2]);} \
\r
195 template<> struct VecTraits<type ## 4> \
\r
197 typedef type elem_type; \
\r
199 static __device__ __host__ __forceinline__ type ## 4 all(type v) {return make_ ## type ## 4(v, v, v, v);} \
\r
200 static __device__ __host__ __forceinline__ type ## 4 make(type x, type y, type z, type w) {return make_ ## type ## 4(x, y, z, w);} \
\r
201 static __device__ __host__ __forceinline__ type ## 4 make(const type* v) {return make_ ## type ## 4(v[0], v[1], v[2], v[3]);} \
\r
203 template<> struct VecTraits<type ## 8> \
\r
205 typedef type elem_type; \
\r
207 static __device__ __host__ __forceinline__ type ## 8 all(type v) {return make_ ## type ## 8(v, v, v, v, v, v, v, v);} \
\r
208 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);} \
\r
209 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]);} \
\r
212 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uchar)
\r
213 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(ushort)
\r
214 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(short)
\r
215 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(int)
\r
216 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uint)
\r
217 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(float)
\r
218 OPENCV_GPU_IMPLEMENT_VEC_TRAITS(double)
\r
220 #undef OPENCV_GPU_IMPLEMENT_VEC_TRAITS
\r
222 template<> struct VecTraits<char>
\r
224 typedef char elem_type;
\r
226 static __device__ __host__ __forceinline__ char all(char v) {return v;}
\r
227 static __device__ __host__ __forceinline__ char make(char x) {return x;}
\r
228 static __device__ __host__ __forceinline__ char make(const char* x) {return *x;}
\r
230 template<> struct VecTraits<schar>
\r
232 typedef schar elem_type;
\r
234 static __device__ __host__ __forceinline__ schar all(schar v) {return v;}
\r
235 static __device__ __host__ __forceinline__ schar make(schar x) {return x;}
\r
236 static __device__ __host__ __forceinline__ schar make(const schar* x) {return *x;}
\r
238 template<> struct VecTraits<char1>
\r
240 typedef schar elem_type;
\r
242 static __device__ __host__ __forceinline__ char1 all(schar v) {return make_char1(v);}
\r
243 static __device__ __host__ __forceinline__ char1 make(schar x) {return make_char1(x);}
\r
244 static __device__ __host__ __forceinline__ char1 make(const schar* v) {return make_char1(v[0]);}
\r
246 template<> struct VecTraits<char2>
\r
248 typedef schar elem_type;
\r
250 static __device__ __host__ __forceinline__ char2 all(schar v) {return make_char2(v, v);}
\r
251 static __device__ __host__ __forceinline__ char2 make(schar x, schar y) {return make_char2(x, y);}
\r
252 static __device__ __host__ __forceinline__ char2 make(const schar* v) {return make_char2(v[0], v[1]);}
\r
254 template<> struct VecTraits<char3>
\r
256 typedef schar elem_type;
\r
258 static __device__ __host__ __forceinline__ char3 all(schar v) {return make_char3(v, v, v);}
\r
259 static __device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z) {return make_char3(x, y, z);}
\r
260 static __device__ __host__ __forceinline__ char3 make(const schar* v) {return make_char3(v[0], v[1], v[2]);}
\r
262 template<> struct VecTraits<char4>
\r
264 typedef schar elem_type;
\r
266 static __device__ __host__ __forceinline__ char4 all(schar v) {return make_char4(v, v, v, v);}
\r
267 static __device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w) {return make_char4(x, y, z, w);}
\r
268 static __device__ __host__ __forceinline__ char4 make(const schar* v) {return make_char4(v[0], v[1], v[2], v[3]);}
\r
270 template<> struct VecTraits<char8>
\r
272 typedef schar elem_type;
\r
274 static __device__ __host__ __forceinline__ char8 all(schar v) {return make_char8(v, v, v, v, v, v, v, v);}
\r
275 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);}
\r
276 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]);}
\r
278 }}} // namespace cv { namespace gpu { namespace device
\r
280 #endif // __OPENCV_GPU_VEC_TRAITS_HPP__
\r