vec_traits.hpp
Go to the documentation of this file.
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_GPU_VEC_TRAITS_HPP__
44 #define __OPENCV_GPU_VEC_TRAITS_HPP__
45 
46 #include "common.hpp"
47 
48 namespace cv { namespace gpu { namespace device
49 {
50  template<typename T, int N> struct TypeVec;
51 
52  struct __align__(8) uchar8
53  {
54  uchar a0, a1, a2, a3, a4, a5, a6, a7;
55  };
56  static __host__ __device__ __forceinline__ uchar8 make_uchar8(uchar a0, uchar a1, uchar a2, uchar a3, uchar a4, uchar a5, uchar a6, uchar a7)
57  {
58  uchar8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
59  return val;
60  }
61  struct __align__(8) char8
62  {
63  schar a0, a1, a2, a3, a4, a5, a6, a7;
64  };
65  static __host__ __device__ __forceinline__ char8 make_char8(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7)
66  {
67  char8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
68  return val;
69  }
70  struct __align__(16) ushort8
71  {
72  ushort a0, a1, a2, a3, a4, a5, a6, a7;
73  };
74  static __host__ __device__ __forceinline__ ushort8 make_ushort8(ushort a0, ushort a1, ushort a2, ushort a3, ushort a4, ushort a5, ushort a6, ushort a7)
75  {
76  ushort8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
77  return val;
78  }
79  struct __align__(16) short8
80  {
81  short a0, a1, a2, a3, a4, a5, a6, a7;
82  };
83  static __host__ __device__ __forceinline__ short8 make_short8(short a0, short a1, short a2, short a3, short a4, short a5, short a6, short a7)
84  {
85  short8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
86  return val;
87  }
88  struct __align__(32) uint8
89  {
90  uint a0, a1, a2, a3, a4, a5, a6, a7;
91  };
92  static __host__ __device__ __forceinline__ uint8 make_uint8(uint a0, uint a1, uint a2, uint a3, uint a4, uint a5, uint a6, uint a7)
93  {
94  uint8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
95  return val;
96  }
97  struct __align__(32) int8
98  {
99  int a0, a1, a2, a3, a4, a5, a6, a7;
100  };
101  static __host__ __device__ __forceinline__ int8 make_int8(int a0, int a1, int a2, int a3, int a4, int a5, int a6, int a7)
102  {
103  int8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
104  return val;
105  }
106  struct __align__(32) float8
107  {
108  float a0, a1, a2, a3, a4, a5, a6, a7;
109  };
110  static __host__ __device__ __forceinline__ float8 make_float8(float a0, float a1, float a2, float a3, float a4, float a5, float a6, float a7)
111  {
112  float8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
113  return val;
114  }
115  struct double8
116  {
117  double a0, a1, a2, a3, a4, a5, a6, a7;
118  };
119  static __host__ __device__ __forceinline__ double8 make_double8(double a0, double a1, double a2, double a3, double a4, double a5, double a6, double a7)
120  {
121  double8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
122  return val;
123  }
124 
125 #define OPENCV_GPU_IMPLEMENT_TYPE_VEC(type) \
126  template<> struct TypeVec<type, 1> { typedef type vec_type; }; \
127  template<> struct TypeVec<type ## 1, 1> { typedef type ## 1 vec_type; }; \
128  template<> struct TypeVec<type, 2> { typedef type ## 2 vec_type; }; \
129  template<> struct TypeVec<type ## 2, 2> { typedef type ## 2 vec_type; }; \
130  template<> struct TypeVec<type, 3> { typedef type ## 3 vec_type; }; \
131  template<> struct TypeVec<type ## 3, 3> { typedef type ## 3 vec_type; }; \
132  template<> struct TypeVec<type, 4> { typedef type ## 4 vec_type; }; \
133  template<> struct TypeVec<type ## 4, 4> { typedef type ## 4 vec_type; }; \
134  template<> struct TypeVec<type, 8> { typedef type ## 8 vec_type; }; \
135  template<> struct TypeVec<type ## 8, 8> { typedef type ## 8 vec_type; };
136 
145 
146  #undef OPENCV_GPU_IMPLEMENT_TYPE_VEC
147 
148  template<> struct TypeVec<schar, 1> { typedef schar vec_type; };
149  template<> struct TypeVec<schar, 2> { typedef char2 vec_type; };
150  template<> struct TypeVec<schar, 3> { typedef char3 vec_type; };
151  template<> struct TypeVec<schar, 4> { typedef char4 vec_type; };
152  template<> struct TypeVec<schar, 8> { typedef char8 vec_type; };
153 
154  template<> struct TypeVec<bool, 1> { typedef uchar vec_type; };
155  template<> struct TypeVec<bool, 2> { typedef uchar2 vec_type; };
156  template<> struct TypeVec<bool, 3> { typedef uchar3 vec_type; };
157  template<> struct TypeVec<bool, 4> { typedef uchar4 vec_type; };
158  template<> struct TypeVec<bool, 8> { typedef uchar8 vec_type; };
159 
160  template<typename T> struct VecTraits;
161 
162 #define OPENCV_GPU_IMPLEMENT_VEC_TRAITS(type) \
163  template<> struct VecTraits<type> \
164  { \
165  typedef type elem_type; \
166  enum {cn=1}; \
167  static __device__ __host__ __forceinline__ type all(type v) {return v;} \
168  static __device__ __host__ __forceinline__ type make(type x) {return x;} \
169  static __device__ __host__ __forceinline__ type make(const type* v) {return *v;} \
170  }; \
171  template<> struct VecTraits<type ## 1> \
172  { \
173  typedef type elem_type; \
174  enum {cn=1}; \
175  static __device__ __host__ __forceinline__ type ## 1 all(type v) {return make_ ## type ## 1(v);} \
176  static __device__ __host__ __forceinline__ type ## 1 make(type x) {return make_ ## type ## 1(x);} \
177  static __device__ __host__ __forceinline__ type ## 1 make(const type* v) {return make_ ## type ## 1(*v);} \
178  }; \
179  template<> struct VecTraits<type ## 2> \
180  { \
181  typedef type elem_type; \
182  enum {cn=2}; \
183  static __device__ __host__ __forceinline__ type ## 2 all(type v) {return make_ ## type ## 2(v, v);} \
184  static __device__ __host__ __forceinline__ type ## 2 make(type x, type y) {return make_ ## type ## 2(x, y);} \
185  static __device__ __host__ __forceinline__ type ## 2 make(const type* v) {return make_ ## type ## 2(v[0], v[1]);} \
186  }; \
187  template<> struct VecTraits<type ## 3> \
188  { \
189  typedef type elem_type; \
190  enum {cn=3}; \
191  static __device__ __host__ __forceinline__ type ## 3 all(type v) {return make_ ## type ## 3(v, v, v);} \
192  static __device__ __host__ __forceinline__ type ## 3 make(type x, type y, type z) {return make_ ## type ## 3(x, y, z);} \
193  static __device__ __host__ __forceinline__ type ## 3 make(const type* v) {return make_ ## type ## 3(v[0], v[1], v[2]);} \
194  }; \
195  template<> struct VecTraits<type ## 4> \
196  { \
197  typedef type elem_type; \
198  enum {cn=4}; \
199  static __device__ __host__ __forceinline__ type ## 4 all(type v) {return make_ ## type ## 4(v, v, v, v);} \
200  static __device__ __host__ __forceinline__ type ## 4 make(type x, type y, type z, type w) {return make_ ## type ## 4(x, y, z, w);} \
201  static __device__ __host__ __forceinline__ type ## 4 make(const type* v) {return make_ ## type ## 4(v[0], v[1], v[2], v[3]);} \
202  }; \
203  template<> struct VecTraits<type ## 8> \
204  { \
205  typedef type elem_type; \
206  enum {cn=8}; \
207  static __device__ __host__ __forceinline__ type ## 8 all(type v) {return make_ ## type ## 8(v, v, v, v, v, v, v, v);} \
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);} \
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]);} \
210  };
211 
219 
220  #undef OPENCV_GPU_IMPLEMENT_VEC_TRAITS
221 
222  template<> struct VecTraits<char>
223  {
224  typedef char elem_type;
225  enum {cn=1};
226  static __device__ __host__ __forceinline__ char all(char v) {return v;}
227  static __device__ __host__ __forceinline__ char make(char x) {return x;}
228  static __device__ __host__ __forceinline__ char make(const char* x) {return *x;}
229  };
230  template<> struct VecTraits<schar>
231  {
232  typedef schar elem_type;
233  enum {cn=1};
234  static __device__ __host__ __forceinline__ schar all(schar v) {return v;}
235  static __device__ __host__ __forceinline__ schar make(schar x) {return x;}
236  static __device__ __host__ __forceinline__ schar make(const schar* x) {return *x;}
237  };
238  template<> struct VecTraits<char1>
239  {
240  typedef schar elem_type;
241  enum {cn=1};
242  static __device__ __host__ __forceinline__ char1 all(schar v) {return make_char1(v);}
243  static __device__ __host__ __forceinline__ char1 make(schar x) {return make_char1(x);}
244  static __device__ __host__ __forceinline__ char1 make(const schar* v) {return make_char1(v[0]);}
245  };
246  template<> struct VecTraits<char2>
247  {
248  typedef schar elem_type;
249  enum {cn=2};
250  static __device__ __host__ __forceinline__ char2 all(schar v) {return make_char2(v, v);}
251  static __device__ __host__ __forceinline__ char2 make(schar x, schar y) {return make_char2(x, y);}
252  static __device__ __host__ __forceinline__ char2 make(const schar* v) {return make_char2(v[0], v[1]);}
253  };
254  template<> struct VecTraits<char3>
255  {
256  typedef schar elem_type;
257  enum {cn=3};
258  static __device__ __host__ __forceinline__ char3 all(schar v) {return make_char3(v, v, v);}
259  static __device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z) {return make_char3(x, y, z);}
260  static __device__ __host__ __forceinline__ char3 make(const schar* v) {return make_char3(v[0], v[1], v[2]);}
261  };
262  template<> struct VecTraits<char4>
263  {
264  typedef schar elem_type;
265  enum {cn=4};
266  static __device__ __host__ __forceinline__ char4 all(schar v) {return make_char4(v, v, v, v);}
267  static __device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w) {return make_char4(x, y, z, w);}
268  static __device__ __host__ __forceinline__ char4 make(const schar* v) {return make_char4(v[0], v[1], v[2], v[3]);}
269  };
270  template<> struct VecTraits<char8>
271  {
272  typedef schar elem_type;
273  enum {cn=8};
274  static __device__ __host__ __forceinline__ char8 all(schar v) {return make_char8(v, v, v, v, v, v, v, v);}
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);}
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]);}
277  };
278 }}} // namespace cv { namespace gpu { namespace device
279 
280 #endif // __OPENCV_GPU_VEC_TRAITS_HPP__
uchar3 vec_type
Definition: vec_traits.hpp:156
short float uchar uchar uchar uchar uchar ushort int uchar ushort int float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float int int int float int int int float int CV_CUDEV_IMPLEMENT_VEC_BINARY_OP char CV_CUDEV_IMPLEMENT_VEC_BINARY_OP ushort CV_CUDEV_IMPLEMENT_VEC_BINARY_OP short CV_CUDEV_IMPLEMENT_VEC_BINARY_OP int CV_CUDEV_IMPLEMENT_VEC_BINARY_OP uint CV_CUDEV_IMPLEMENT_VEC_BINARY_OP float CV_CUDEV_IMPLEMENT_VEC_BINARY_OP double char
Definition: vec_math.hpp:426
GLenum GLint GLint y
Definition: core_c.h:613
schar vec_type
Definition: vec_traits.hpp:148
schar elem_type
Definition: vec_traits.hpp:240
signed char schar
Definition: types_c.h:174
schar elem_type
Definition: vec_traits.hpp:264
static __device__ __host__ __forceinline__ char4 all(schar v)
Definition: vec_traits.hpp:266
static __device__ __host__ __forceinline__ schar make(schar x)
Definition: vec_traits.hpp:235
static __device__ __host__ __forceinline__ char3 all(schar v)
Definition: vec_traits.hpp:258
static __device__ __host__ __forceinline__ char8 make(schar a0, schar a1, schar a2, schar a3, schar a4, schar a5, schar a6, schar a7)
Definition: vec_traits.hpp:275
static __device__ __host__ __forceinline__ schar make(const schar *x)
Definition: vec_traits.hpp:236
uchar vec_type
Definition: vec_traits.hpp:154
double a5
Definition: vec_traits.hpp:117
static __device__ __host__ __forceinline__ char1 make(const schar *v)
Definition: vec_traits.hpp:244
double a4
Definition: vec_traits.hpp:117
static __device__ __host__ __forceinline__ char2 make(schar x, schar y)
Definition: vec_traits.hpp:251
char2 vec_type
Definition: vec_traits.hpp:149
uchar8 vec_type
Definition: vec_traits.hpp:158
double a2
Definition: vec_traits.hpp:117
schar elem_type
Definition: vec_traits.hpp:272
struct __align__(8) uchar8
Definition: vec_traits.hpp:52
signed char schar
Definition: common.hpp:102
unsigned char uchar
Definition: common.hpp:100
double a0
Definition: vec_traits.hpp:117
GLuint GLfloat * val
double a6
Definition: vec_traits.hpp:117
schar elem_type
Definition: vec_traits.hpp:256
double a3
Definition: vec_traits.hpp:117
static __device__ __host__ __forceinline__ char2 make(const schar *v)
Definition: vec_traits.hpp:252
static __device__ __host__ __forceinline__ char8 make(const schar *v)
Definition: vec_traits.hpp:276
static __device__ __host__ __forceinline__ char1 make(schar x)
Definition: vec_traits.hpp:243
GLenum GLint x
Definition: core_c.h:632
static __device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w)
Definition: vec_traits.hpp:267
unsigned int uint
Definition: common.hpp:104
static __device__ __host__ __forceinline__ schar all(schar v)
Definition: vec_traits.hpp:234
schar elem_type
Definition: vec_traits.hpp:248
const GLdouble * v
unsigned short ushort
Definition: common.hpp:101
static __device__ __host__ __forceinline__ char4 make(const schar *v)
Definition: vec_traits.hpp:268
GLdouble GLdouble z
char4 vec_type
Definition: vec_traits.hpp:151
unsigned short ushort
Definition: types_c.h:171
static __device__ __host__ __forceinline__ char2 all(schar v)
Definition: vec_traits.hpp:250
char8 vec_type
Definition: vec_traits.hpp:152
uchar2 vec_type
Definition: vec_traits.hpp:155
OPENCV_GPU_IMPLEMENT_TYPE_VEC(uchar) OPENCV_GPU_IMPLEMENT_TYPE_VEC(char) OPENCV_GPU_IMPLEMENT_TYPE_VEC(ushort) OPENCV_GPU_IMPLEMENT_TYPE_VEC(short) OPENCV_GPU_IMPLEMENT_TYPE_VEC(int) OPENCV_GPU_IMPLEMENT_TYPE_VEC(uint) OPENCV_GPU_IMPLEMENT_TYPE_VEC(float) OPENCV_GPU_IMPLEMENT_TYPE_VEC(double) template<> struct TypeVec< schar
Definition: vec_traits.hpp:50
unsigned char uchar
Definition: types_c.h:170
schar elem_type
Definition: vec_traits.hpp:232
OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uchar) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(ushort) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(short) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(int) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(uint) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(float) OPENCV_GPU_IMPLEMENT_VEC_TRAITS(double) template<> struct VecTraits< char >
Definition: vec_traits.hpp:212
GLubyte GLubyte GLubyte GLubyte w
double a7
Definition: vec_traits.hpp:117
static __device__ __host__ __forceinline__ char8 all(schar v)
Definition: vec_traits.hpp:274
int x
Definition: highgui_c.h:186
double a1
Definition: vec_traits.hpp:117
static __device__ __host__ __forceinline__ char1 all(schar v)
Definition: vec_traits.hpp:242
uchar4 vec_type
Definition: vec_traits.hpp:157
static __device__ __host__ __forceinline__ char3 make(const schar *v)
Definition: vec_traits.hpp:260
Definition: vec_traits.hpp:160
char3 vec_type
Definition: vec_traits.hpp:150
Definition: vec_traits.hpp:115
static __device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z)
Definition: vec_traits.hpp:259