OpenCV  2.4.13.7
Open Source Computer Vision
Classes | Functions
cv::gpu::device::reduce_key_val_detail Namespace Reference

Classes

struct  Dispatcher
 
struct  For
 
struct  For< N, N >
 
struct  Generic
 
struct  GenericOptimized32
 
struct  GetType
 
struct  GetType< T & >
 
struct  GetType< T * >
 
struct  GetType< volatile T * >
 
struct  IsPowerOf2
 
struct  StaticIf
 
struct  StaticIf< false, T1, T2 >
 
struct  StaticIf< true, T1, T2 >
 
struct  Unroll
 
struct  Unroll< 0, KP, KR, VP, VR, Cmp >
 
struct  WarpOptimized
 

Functions

template<typename T >
__device__ __forceinline__ void loadToSmem (volatile T *smem, T &data, unsigned int tid)
 
template<typename T >
__device__ __forceinline__ void loadFromSmem (volatile T *smem, T &data, unsigned int tid)
 
template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void loadToSmem (const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &smem, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &data, unsigned int tid)
 
template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void loadFromSmem (const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &smem, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &data, unsigned int tid)
 
template<typename V >
__device__ __forceinline__ void copyValsShfl (V &val, unsigned int delta, int width)
 
template<typename V >
__device__ __forceinline__ void copyVals (volatile V *svals, V &val, unsigned int tid, unsigned int delta)
 
template<typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void copyValsShfl (const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, unsigned int delta, int width)
 
template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void copyVals (const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &svals, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, unsigned int tid, unsigned int delta)
 
template<typename K , typename V , class Cmp >
__device__ __forceinline__ void mergeShfl (K &key, V &val, const Cmp &cmp, unsigned int delta, int width)
 
template<typename K , typename V , class Cmp >
__device__ __forceinline__ void merge (volatile K *skeys, K &key, volatile V *svals, V &val, const Cmp &cmp, unsigned int tid, unsigned int delta)
 
template<typename K , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp >
__device__ __forceinline__ void mergeShfl (K &key, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, const Cmp &cmp, unsigned int delta, int width)
 
template<typename K , typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp >
__device__ __forceinline__ void merge (volatile K *skeys, K &key, const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &svals, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, const Cmp &cmp, unsigned int tid, unsigned int delta)
 
template<typename KR0 , typename KR1 , typename KR2 , typename KR3 , typename KR4 , typename KR5 , typename KR6 , typename KR7 , typename KR8 , typename KR9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp0 , class Cmp1 , class Cmp2 , class Cmp3 , class Cmp4 , class Cmp5 , class Cmp6 , class Cmp7 , class Cmp8 , class Cmp9 >
__device__ __forceinline__ void mergeShfl (const thrust::tuple< KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9 > &key, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, const thrust::tuple< Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9 > &cmp, unsigned int delta, int width)
 
template<typename KP0 , typename KP1 , typename KP2 , typename KP3 , typename KP4 , typename KP5 , typename KP6 , typename KP7 , typename KP8 , typename KP9 , typename KR0 , typename KR1 , typename KR2 , typename KR3 , typename KR4 , typename KR5 , typename KR6 , typename KR7 , typename KR8 , typename KR9 , typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp0 , class Cmp1 , class Cmp2 , class Cmp3 , class Cmp4 , class Cmp5 , class Cmp6 , class Cmp7 , class Cmp8 , class Cmp9 >
__device__ __forceinline__ void merge (const thrust::tuple< KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9 > &skeys, const thrust::tuple< KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9 > &key, const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &svals, const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &val, const thrust::tuple< Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9 > &cmp, unsigned int tid, unsigned int delta)
 

Function Documentation

§ copyVals() [1/2]

template<typename V >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::copyVals ( volatile V *  svals,
V &  val,
unsigned int  tid,
unsigned int  delta 
)

§ copyVals() [2/2]

template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::copyVals ( const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &  svals,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
unsigned int  tid,
unsigned int  delta 
)

§ copyValsShfl() [1/2]

template<typename V >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::copyValsShfl ( V &  val,
unsigned int  delta,
int  width 
)

§ copyValsShfl() [2/2]

template<typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::copyValsShfl ( const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
unsigned int  delta,
int  width 
)

§ loadFromSmem() [1/2]

template<typename T >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::loadFromSmem ( volatile T *  smem,
T &  data,
unsigned int  tid 
)

§ loadFromSmem() [2/2]

template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::loadFromSmem ( const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &  smem,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  data,
unsigned int  tid 
)

§ loadToSmem() [1/2]

template<typename T >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::loadToSmem ( volatile T *  smem,
T &  data,
unsigned int  tid 
)

§ loadToSmem() [2/2]

template<typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::loadToSmem ( const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &  smem,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  data,
unsigned int  tid 
)

§ merge() [1/3]

template<typename K , typename V , class Cmp >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::merge ( volatile K *  skeys,
K &  key,
volatile V *  svals,
V &  val,
const Cmp &  cmp,
unsigned int  tid,
unsigned int  delta 
)

§ merge() [2/3]

template<typename K , typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::merge ( volatile K *  skeys,
K &  key,
const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &  svals,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
const Cmp &  cmp,
unsigned int  tid,
unsigned int  delta 
)

§ merge() [3/3]

template<typename KP0 , typename KP1 , typename KP2 , typename KP3 , typename KP4 , typename KP5 , typename KP6 , typename KP7 , typename KP8 , typename KP9 , typename KR0 , typename KR1 , typename KR2 , typename KR3 , typename KR4 , typename KR5 , typename KR6 , typename KR7 , typename KR8 , typename KR9 , typename VP0 , typename VP1 , typename VP2 , typename VP3 , typename VP4 , typename VP5 , typename VP6 , typename VP7 , typename VP8 , typename VP9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp0 , class Cmp1 , class Cmp2 , class Cmp3 , class Cmp4 , class Cmp5 , class Cmp6 , class Cmp7 , class Cmp8 , class Cmp9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::merge ( const thrust::tuple< KP0, KP1, KP2, KP3, KP4, KP5, KP6, KP7, KP8, KP9 > &  skeys,
const thrust::tuple< KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9 > &  key,
const thrust::tuple< VP0, VP1, VP2, VP3, VP4, VP5, VP6, VP7, VP8, VP9 > &  svals,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
const thrust::tuple< Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9 > &  cmp,
unsigned int  tid,
unsigned int  delta 
)

§ mergeShfl() [1/3]

template<typename K , typename V , class Cmp >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::mergeShfl ( K &  key,
V &  val,
const Cmp &  cmp,
unsigned int  delta,
int  width 
)

§ mergeShfl() [2/3]

template<typename K , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::mergeShfl ( K &  key,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
const Cmp &  cmp,
unsigned int  delta,
int  width 
)

§ mergeShfl() [3/3]

template<typename KR0 , typename KR1 , typename KR2 , typename KR3 , typename KR4 , typename KR5 , typename KR6 , typename KR7 , typename KR8 , typename KR9 , typename VR0 , typename VR1 , typename VR2 , typename VR3 , typename VR4 , typename VR5 , typename VR6 , typename VR7 , typename VR8 , typename VR9 , class Cmp0 , class Cmp1 , class Cmp2 , class Cmp3 , class Cmp4 , class Cmp5 , class Cmp6 , class Cmp7 , class Cmp8 , class Cmp9 >
__device__ __forceinline__ void cv::gpu::device::reduce_key_val_detail::mergeShfl ( const thrust::tuple< KR0, KR1, KR2, KR3, KR4, KR5, KR6, KR7, KR8, KR9 > &  key,
const thrust::tuple< VR0, VR1, VR2, VR3, VR4, VR5, VR6, VR7, VR8, VR9 > &  val,
const thrust::tuple< Cmp0, Cmp1, Cmp2, Cmp3, Cmp4, Cmp5, Cmp6, Cmp7, Cmp8, Cmp9 > &  cmp,
unsigned int  delta,
int  width 
)