43 #ifndef __OPENCV_GPU_WARP_SHUFFLE_HPP__
44 #define __OPENCV_GPU_WARP_SHUFFLE_HPP__
46 namespace cv {
namespace gpu {
namespace device
51 #if __CUDA_ARCH__ >= 300
52 return __shfl(val, srcLane,
width);
59 #if __CUDA_ARCH__ >= 300
60 return (
unsigned int) __shfl((
int) val, srcLane,
width);
67 #if __CUDA_ARCH__ >= 300
68 int lo = __double2loint(val);
69 int hi = __double2hiint(val);
71 lo = __shfl(lo, srcLane,
width);
72 hi = __shfl(hi, srcLane,
width);
74 return __hiloint2double(hi, lo);
83 #if __CUDA_ARCH__ >= 300
84 return __shfl_down(val, delta,
width);
91 #if __CUDA_ARCH__ >= 300
92 return (
unsigned int) __shfl_down((
int) val, delta,
width);
99 #if __CUDA_ARCH__ >= 300
100 int lo = __double2loint(val);
101 int hi = __double2hiint(val);
103 lo = __shfl_down(lo, delta,
width);
104 hi = __shfl_down(hi, delta,
width);
106 return __hiloint2double(hi, lo);
112 template <
typename T>
115 #if __CUDA_ARCH__ >= 300
116 return __shfl_up(val, delta,
width);
123 #if __CUDA_ARCH__ >= 300
124 return (
unsigned int) __shfl_up((
int) val, delta,
width);
131 #if __CUDA_ARCH__ >= 300
132 int lo = __double2loint(val);
133 int hi = __double2hiint(val);
135 lo = __shfl_up(lo, delta,
width);
136 hi = __shfl_up(hi, delta,
width);
138 return __hiloint2double(hi, lo);
145 #endif // __OPENCV_GPU_WARP_SHUFFLE_HPP__
__device__ __forceinline__ T shfl(T val, int srcLane, int width=warpSize)
Definition: warp_shuffle.hpp:49
__device__ __forceinline__ T shfl_down(T val, unsigned int delta, int width=warpSize)
Definition: warp_shuffle.hpp:81
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
__device__ __forceinline__ T shfl_up(T val, unsigned int delta, int width=warpSize)
Definition: warp_shuffle.hpp:113
CvSize int int int CvPoint int delta
Definition: core_c.h:1427
double double double double double double CvSize * warpSize
Definition: legacy.hpp:680