43
#ifndef OPENCV_CUDA_VEC_TRAITS_HPP
 
44
#define OPENCV_CUDA_VEC_TRAITS_HPP
 
54
namespace
cv
{
namespace
cuda {
namespace
device
 
56
template<
typename
T,
int
N>
struct
TypeVec;
 
58
struct
__align__(8) uchar8
 
60
uchar a0, a1, a2, a3, a4, a5, a6, 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)
 
64
uchar8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
67
struct
__align__(8) char8
 
69
schar a0, a1, a2, a3, a4, a5, a6, 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)
 
73
char8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
76
struct
__align__(16) ushort8
 
78
ushort a0, a1, a2, a3, a4, a5, a6, 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)
 
82
ushort8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
85
struct
__align__(16) short8
 
87
short
a0, a1, a2, a3, a4, a5, a6, 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)
 
91
short8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
94
struct
__align__(32) uint8
 
96
uint a0, a1, a2, a3, a4, a5, a6, 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)
 
100
uint8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
103
struct
__align__(32) int8
 
105
int
a0, a1, a2, a3, a4, a5, a6, 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)
 
109
int8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
112
struct
__align__(32) float8
 
114
float
a0, a1, a2, a3, a4, a5, a6, 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)
 
118
float8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
123
double
a0, a1, a2, a3, a4, a5, a6, 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)
 
127
double8 val = {a0, a1, a2, a3, a4, a5, a6, a7};
 
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; };
 
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)
 
152
#undef OPENCV_CUDA_IMPLEMENT_TYPE_VEC
 
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; };
 
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; };
 
166
template<
typename
T>
struct
VecTraits;
 
168
#define OPENCV_CUDA_IMPLEMENT_VEC_TRAITS(type) \
 
169
template<> struct VecTraits<type> \
 
171
typedef type elem_type; \
 
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;} \
 
177
template<> struct VecTraits<type ## 1> \
 
179
typedef type elem_type; \
 
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);} \
 
185
template<> struct VecTraits<type ## 2> \
 
187
typedef type elem_type; \
 
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]);} \
 
193
template<> struct VecTraits<type ## 3> \
 
195
typedef type elem_type; \
 
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]);} \
 
201
template<> struct VecTraits<type ## 4> \
 
203
typedef type elem_type; \
 
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]);} \
 
209
template<> struct VecTraits<type ## 8> \
 
211
typedef type elem_type; \
 
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]);} \
 
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)
 
226
#undef OPENCV_CUDA_IMPLEMENT_VEC_TRAITS
 
228
template<>
struct
VecTraits<char>
 
230
typedef
char
elem_type;
 
232
static
__device__ __host__ __forceinline__
char
all(
char
v) {
return
v;}
 
233
static
__device__ __host__ __forceinline__
char
make(
char
x) {
return
x;}
 
234
static
__device__ __host__ __forceinline__
char
make(
const
char* x) {
return
*x;}
 
236
template<>
struct
VecTraits<schar>
 
238
typedef
schar elem_type;
 
240
static
__device__ __host__ __forceinline__ schar all(schar v) {
return
v;}
 
241
static
__device__ __host__ __forceinline__ schar make(schar x) {
return
x;}
 
242
static
__device__ __host__ __forceinline__ schar make(
const
schar* x) {
return
*x;}
 
244
template<>
struct
VecTraits<char1>
 
246
typedef
schar elem_type;
 
248
static
__device__ __host__ __forceinline__ char1 all(schar v) {
return
make_char1(v);}
 
249
static
__device__ __host__ __forceinline__ char1 make(schar x) {
return
make_char1(x);}
 
250
static
__device__ __host__ __forceinline__ char1 make(
const
schar* v) {
return
make_char1(v[0]);}
 
252
template<>
struct
VecTraits<char2>
 
254
typedef
schar elem_type;
 
256
static
__device__ __host__ __forceinline__ char2 all(schar v) {
return
make_char2(v, v);}
 
257
static
__device__ __host__ __forceinline__ char2 make(schar x, schar y) {
return
make_char2(x, y);}
 
258
static
__device__ __host__ __forceinline__ char2 make(
const
schar* v) {
return
make_char2(v[0], v[1]);}
 
260
template<>
struct
VecTraits<char3>
 
262
typedef
schar elem_type;
 
264
static
__device__ __host__ __forceinline__ char3 all(schar v) {
return
make_char3(v, v, v);}
 
265
static
__device__ __host__ __forceinline__ char3 make(schar x, schar y, schar z) {
return
make_char3(x, y, z);}
 
266
static
__device__ __host__ __forceinline__ char3 make(
const
schar* v) {
return
make_char3(v[0], v[1], v[2]);}
 
268
template<>
struct
VecTraits<char4>
 
270
typedef
schar elem_type;
 
272
static
__device__ __host__ __forceinline__ char4 all(schar v) {
return
make_char4(v, v, v, v);}
 
273
static
__device__ __host__ __forceinline__ char4 make(schar x, schar y, schar z, schar w) {
return
make_char4(x, y, z, w);}
 
274
static
__device__ __host__ __forceinline__ char4 make(
const
schar* v) {
return
make_char4(v[0], v[1], v[2], v[3]);}
 
276
template<>
struct
VecTraits<char8>
 
278
typedef
schar elem_type;
 
280
static
__device__ __host__ __forceinline__ char8 all(schar v) {
return
make_char8(v, v, v, v, v, v, v, v);}
 
281
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);}
 
282
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]);}
 
"black box" representation of the file storage associated with a file on disk.
Definition:
aruco.hpp:75