43 #ifndef __OPENCV_GPU_REDUCE_HPP__
44 #define __OPENCV_GPU_REDUCE_HPP__
46 #include <thrust/tuple.h>
50 namespace cv {
namespace gpu {
namespace device
52 template <
int N,
typename T,
class Op>
53 __device__ __forceinline__
void reduce(
volatile T* smem,
T&
val,
unsigned int tid,
const Op& op)
58 typename P0,
typename P1,
typename P2,
typename P3,
typename P4,
typename P5,
typename P6,
typename P7,
typename P8,
typename P9,
59 typename R0,
typename R1,
typename R2,
typename R3,
typename R4,
typename R5,
typename R6,
typename R7,
typename R8,
typename R9,
60 class Op0,
class Op1,
class Op2,
class Op3,
class Op4,
class Op5,
class Op6,
class Op7,
class Op8,
class Op9>
61 __device__ __forceinline__
void reduce(
const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>& smem,
62 const thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>&
val,
64 const thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>& op)
67 const thrust::tuple<P0, P1, P2, P3, P4, P5, P6, P7, P8, P9>&,
68 const thrust::tuple<R0, R1, R2, R3, R4, R5, R6, R7, R8, R9>&,
69 const thrust::tuple<Op0, Op1, Op2, Op3, Op4, Op5, Op6, Op7, Op8, Op9>&>(smem, val, tid, op);
72 template <
unsigned int N,
typename K,
typename V,
class Cmp>
73 __device__ __forceinline__
void reduceKeyVal(
volatile K* skeys, K&
key,
volatile V* svals,
V&
val,
unsigned int tid,
const Cmp& cmp)
77 template <
unsigned int N,
79 typename VP0,
typename VP1,
typename VP2,
typename VP3,
typename VP4,
typename VP5,
typename VP6,
typename VP7,
typename VP8,
typename VP9,
80 typename VR0,
typename VR1,
typename VR2,
typename VR3,
typename VR4,
typename VR5,
typename VR6,
typename VR7,
typename VR8,
typename VR9,
83 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
84 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&
val,
85 unsigned int tid,
const Cmp& cmp)
88 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
89 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
90 const Cmp&>(skeys,
key, svals, val, tid, cmp);
92 template <
unsigned int N,
93 typename KP0,
typename KP1,
typename KP2,
typename KP3,
typename KP4,
typename KP5,
typename KP6,
typename KP7,
typename KP8,
typename KP9,
94 typename KR0,
typename KR1,
typename KR2,
typename KR3,
typename KR4,
typename KR5,
typename KR6,
typename KR7,
typename KR8,
typename KR9,
95 typename VP0,
typename VP1,
typename VP2,
typename VP3,
typename VP4,
typename VP5,
typename VP6,
typename VP7,
typename VP8,
typename VP9,
96 typename VR0,
typename VR1,
typename VR2,
typename VR3,
typename VR4,
typename VR5,
typename VR6,
typename VR7,
typename VR8,
typename VR9,
97 class Cmp0,
class Cmp1,
class Cmp2,
class Cmp3,
class Cmp4,
class Cmp5,
class Cmp6,
class Cmp7,
class Cmp8,
class Cmp9>
98 __device__ __forceinline__
void reduceKeyVal(
const thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>& skeys,
99 const thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>&
key,
100 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>& svals,
101 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&
val,
103 const thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>& cmp)
106 const thrust::tuple<KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9>&,
107 const thrust::tuple<KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9>&,
108 const thrust::tuple<VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9>&,
109 const thrust::tuple<VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9>&,
110 const thrust::tuple<Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9>&
111 >(skeys,
key, svals, val, tid, cmp);
116 template <
typename T0>
117 __device__ __forceinline__
118 thrust::tuple<volatile T0*>
124 template <
typename T0,
typename T1>
125 __device__ __forceinline__
126 thrust::tuple<volatile T0*, volatile T1*>
132 template <
typename T0,
typename T1,
typename T2>
133 __device__ __forceinline__
134 thrust::tuple<volatile T0*, volatile T1*, volatile T2*>
140 template <
typename T0,
typename T1,
typename T2,
typename T3>
141 __device__ __forceinline__
142 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*>
145 return thrust::make_tuple((
volatile T0*) t0, (
volatile T1*) t1, (
volatile T2*) t2, (
volatile T3*) t3);
148 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4>
149 __device__ __forceinline__
150 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*>
153 return thrust::make_tuple((
volatile T0*) t0, (
volatile T1*) t1, (
volatile T2*) t2, (
volatile T3*) t3, (
volatile T4*) t4);
156 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4,
typename T5>
157 __device__ __forceinline__
158 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*>
161 return thrust::make_tuple((
volatile T0*) t0, (
volatile T1*) t1, (
volatile T2*) t2, (
volatile T3*) t3, (
volatile T4*) t4, (
volatile T5*) t5);
164 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4,
typename T5,
typename T6>
165 __device__ __forceinline__
166 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*>
167 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6)
169 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);
172 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4,
typename T5,
typename T6,
typename T7>
173 __device__ __forceinline__
174 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*>
175 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7)
177 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);
180 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4,
typename T5,
typename T6,
typename T7,
typename T8>
181 __device__ __forceinline__
182 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*>
183 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8)
185 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);
188 template <
typename T0,
typename T1,
typename T2,
typename T3,
typename T4,
typename T5,
typename T6,
typename T7,
typename T8,
typename T9>
189 __device__ __forceinline__
190 thrust::tuple<volatile T0*, volatile T1*, volatile T2*, volatile T3*, volatile T4*, volatile T5*, volatile T6*, volatile T7*, volatile T8*, volatile T9*>
191 smem_tuple(T0* t0, T1* t1, T2* t2, T3* t3, T4* t4, T5* t5, T6* t6, T7* t7, T8* t8, T9* t9)
193 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);
197 #endif // __OPENCV_GPU_UTILITY_HPP__
__device__ __forceinline__ void reduce(volatile T *smem, T &val, unsigned int tid, const Op &op)
Definition: reduce.hpp:53
const CvMat const CvMat const CvMat CvSize const CvMat const CvMat CvMat CvMat * R2
Definition: calib3d.hpp:284
CvFileNode const CvStringHashNode * key
Definition: core_c.h:1584
Definition: reduce.hpp:346
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
tuple make_tuple()
Definition: ts_gtest.h:1355
const CvMat const CvMat const CvMat CvSize const CvMat const CvMat CvMat * R1
Definition: calib3d.hpp:284
Definition: reduce_key_val.hpp:483
const CvMat const CvMat const CvMat CvSize const CvMat const CvMat CvMat CvMat CvMat * P1
Definition: calib3d.hpp:284
const CvArr const CvArr * V
Definition: core_c.h:733
__device__ __forceinline__ thrust::tuple< volatile T0 * > smem_tuple(T0 *t0)
Definition: reduce.hpp:119
__device__ __forceinline__ void reduceKeyVal(volatile K *skeys, K &key, volatile V *svals, V &val, unsigned int tid, const Cmp &cmp)
Definition: reduce.hpp:73
const CvMat const CvMat const CvMat CvSize const CvMat const CvMat CvMat CvMat CvMat CvMat * P2
Definition: calib3d.hpp:284