43
#ifndef OPENCV_CUDA_REDUCE_HPP
44
#define OPENCV_CUDA_REDUCE_HPP
47
#define THRUST_DEBUG 0
50
#include <thrust/tuple.h>
51
#include "detail/reduce.hpp"
52
#include "detail/reduce_key_val.hpp"
60
namespace
cv
{
namespace
cuda {
namespace
device
62
template
<
int
N,
typename
T,
class
Op>
63
__device__ __forceinline__
void
reduce(
volatile
T* smem, T& val,
unsigned
int
tid,
const
Op& op)
65
reduce_detail::Dispatcher<N>::reductor::template reduce<volatile T*, T&, const Op&>(smem, val, tid, op);
68
typename
P0,
typename
P1,
typename
P2,
typename
P3,
typename
P4,
typename
P5,
typename
P6,
typename
P7,
typename
P8,
typename
P9,
69
typename
R0,
typename
R1,
typename
R2,
typename
R3,
typename
R4,
typename
R5,
typename
R6,
typename
R7,
typename
R8,
typename
R9,
70
class
Op0,
class
Op1,
class
Op2,
class
Op3,
class
Op4,
class
Op5,
class
Op6,
class
Op7,
class
Op8,
class
Op9>
71
__device__ __forceinline__
void
reduce(
const
thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
72
const
thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>& val,
74
const
thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>& op)
76
reduce_detail::Dispatcher<N>::reductor::template
reduce<
77
const
thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>&,
78
const
thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>&,
79
const
thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>&>(smem, val, tid, op);
82
template
<
unsigned
int
N,
typename
K,
typename
V,
class
Cmp>
83
__device__ __forceinline__
void
reduceKeyVal(
volatile
K* skeys, K& key,
volatile
V* svals, V& val,
unsigned
int
tid,
const
Cmp& cmp)
85
reduce_key_val_detail::Dispatcher<N>::reductor::template reduce<volatile K*, K&, volatile V*, V&, const Cmp&>(skeys, key, svals, val, tid, cmp);
87
template
<
unsigned
int
N,
89
typename
VP0,
typename
VP1,
typename
VP2,
typename
VP3,
typename
VP4,
typename
VP5,
typename
VP6,
typename
VP7,
typename
VP8,
typename
VP9,
90
typename
VR0,
typename
VR1,
typename
VR2,
typename
VR3,
typename
VR4,
typename
VR5,
typename
VR6,
typename
VR7,
typename
VR8,
typename
VR9,
92
__device__ __forceinline__
void
reduceKeyVal(
volatile
K* skeys, K& key,
93
const
thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
94
const
thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
95
unsigned
int
tid,
const
Cmp& cmp)
97
reduce_key_val_detail::Dispatcher<N>::reductor::template
reduce<
volatile
K*, K&,
98
const
thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
99
const
thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
100
const
Cmp&>(skeys, key, svals, val, tid, cmp);
102
template
<
unsigned
int
N,
103
typename
KP0,
typename
KP1,
typename
KP2,
typename
KP3,
typename
KP4,
typename
KP5,
typename
KP6,
typename
KP7,
typename
KP8,
typename
KP9,
104
typename
KR0,
typename
KR1,
typename
KR2,
typename
KR3,
typename
KR4,
typename
KR5,
typename
KR6,
typename
KR7,
typename
KR8,
typename
KR9,
105
typename
VP0,
typename
VP1,
typename
VP2,
typename
VP3,
typename
VP4,
typename
VP5,
typename
VP6,
typename
VP7,
typename
VP8,
typename
VP9,
106
typename
VR0,
typename
VR1,
typename
VR2,
typename
VR3,
typename
VR4,
typename
VR5,
typename
VR6,
typename
VR7,
typename
VR8,
typename
VR9,
107
class
Cmp0,
class
Cmp1,
class
Cmp2,
class
Cmp3,
class
Cmp4,
class
Cmp5,
class
Cmp6,
class
Cmp7,
class
Cmp8,
class
Cmp9>
108
__device__ __forceinline__
void
reduceKeyVal(
const
thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>& skeys,
109
const
thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>& key,
110
const
thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
111
const
thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>& val,
113
const
thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>& cmp)
115
reduce_key_val_detail::Dispatcher<N>::reductor::template
reduce<
116
const
thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>&,
117
const
thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>&,
118
const
thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
119
const
thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
120
const
thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>&
121
>(skeys, key, svals, val, tid, cmp);
126
template
<
typename
T0>
127
__device__ __forceinline__
128
thrust::tuple<volatile T0*>
131
return
thrust::make_tuple((
volatile
T0*) t0);
134
template
<
typename
T0,
typename
T1>
135
__device__ __forceinline__
136
thrust::tuple<volatile T0*, volatile T1*>
137
smem_tuple(T0* t0, T1* t1)
139
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1);
142
template
<
typename
T0,
typename
T1,
typename
T2>
143
__device__ __forceinline__
144
thrust::tuple<volatile T0*, volatile T1*, volatile T2*>
145
smem_tuple(T0* t0, T1* t1, T2* t2)
147
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2);
150
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3>
151
__device__ __forceinline__
152
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*>
153
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3)
155
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3);
158
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4>
159
__device__ __forceinline__
160
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*>
161
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4)
163
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4);
166
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4,
typename
T5>
167
__device__ __forceinline__
168
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*>
169
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5)
171
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4, (
volatile
T5*) t5);
174
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4,
typename
T5,
typename
T6>
175
__device__ __forceinline__
176
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*>
177
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6)
179
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4, (
volatile
T5*) t5, (
volatile
T6*) t6);
182
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4,
typename
T5,
typename
T6,
typename
T7>
183
__device__ __forceinline__
184
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*>
185
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7)
187
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4, (
volatile
T5*) t5, (
volatile
T6*) t6, (
volatile
T7*) t7);
190
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4,
typename
T5,
typename
T6,
typename
T7,
typename
T8>
191
__device__ __forceinline__
192
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*>
193
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8)
195
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4, (
volatile
T5*) t5, (
volatile
T6*) t6, (
volatile
T7*) t7, (
volatile
T8*) t8);
198
template
<
typename
T0,
typename
T1,
typename
T2,
typename
T3,
typename
T4,
typename
T5,
typename
T6,
typename
T7,
typename
T8,
typename
T9>
199
__device__ __forceinline__
200
thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*, volatile T9*>
201
smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8, T9* t9)
203
return
thrust::make_tuple((
volatile
T0*) t0, (
volatile
T1*) t1, (
volatile
T2*) t2, (
volatile
T3*) t3, (
volatile
T4*) t4, (
volatile
T5*) t5, (
volatile
T6*) t6, (
volatile
T7*) t7, (
volatile
T8*) t8, (
volatile
T9*) t9);
CV_EXPORTS_W void reduce(InputArray src, OutputArray dst, int dim, int rtype, int dtype=-1)
Reduces a matrix to a vector.
"black box" representation of the file storage associated with a file on disk.
Definition:
aruco.hpp:75