warp_shuffle.hpp
Go to the documentation of this file.
1 /*M///////////////////////////////////////////////////////////////////////////////////////
2 //
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
4 //
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
8 //
9 //
10 // License Agreement
11 // For Open Source Computer Vision Library
12 //
13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
22 //
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
26 //
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42 
43 #ifndef __OPENCV_GPU_WARP_SHUFFLE_HPP__
44 #define __OPENCV_GPU_WARP_SHUFFLE_HPP__
45 
46 namespace cv { namespace gpu { namespace device
47 {
48  template <typename T>
49  __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize)
50  {
51  #if __CUDA_ARCH__ >= 300
52  return __shfl(val, srcLane, width);
53  #else
54  return T();
55  #endif
56  }
57  __device__ __forceinline__ unsigned int shfl(unsigned int val, int srcLane, int width = warpSize)
58  {
59  #if __CUDA_ARCH__ >= 300
60  return (unsigned int) __shfl((int) val, srcLane, width);
61  #else
62  return 0;
63  #endif
64  }
65  __device__ __forceinline__ double shfl(double val, int srcLane, int width = warpSize)
66  {
67  #if __CUDA_ARCH__ >= 300
68  int lo = __double2loint(val);
69  int hi = __double2hiint(val);
70 
71  lo = __shfl(lo, srcLane, width);
72  hi = __shfl(hi, srcLane, width);
73 
74  return __hiloint2double(hi, lo);
75  #else
76  return 0.0;
77  #endif
78  }
79 
80  template <typename T>
81  __device__ __forceinline__ T shfl_down(T val, unsigned int delta, int width = warpSize)
82  {
83  #if __CUDA_ARCH__ >= 300
84  return __shfl_down(val, delta, width);
85  #else
86  return T();
87  #endif
88  }
89  __device__ __forceinline__ unsigned int shfl_down(unsigned int val, unsigned int delta, int width = warpSize)
90  {
91  #if __CUDA_ARCH__ >= 300
92  return (unsigned int) __shfl_down((int) val, delta, width);
93  #else
94  return 0;
95  #endif
96  }
97  __device__ __forceinline__ double shfl_down(double val, unsigned int delta, int width = warpSize)
98  {
99  #if __CUDA_ARCH__ >= 300
100  int lo = __double2loint(val);
101  int hi = __double2hiint(val);
102 
103  lo = __shfl_down(lo, delta, width);
104  hi = __shfl_down(hi, delta, width);
105 
106  return __hiloint2double(hi, lo);
107  #else
108  return 0.0;
109  #endif
110  }
111 
112  template <typename T>
113  __device__ __forceinline__ T shfl_up(T val, unsigned int delta, int width = warpSize)
114  {
115  #if __CUDA_ARCH__ >= 300
116  return __shfl_up(val, delta, width);
117  #else
118  return T();
119  #endif
120  }
121  __device__ __forceinline__ unsigned int shfl_up(unsigned int val, unsigned int delta, int width = warpSize)
122  {
123  #if __CUDA_ARCH__ >= 300
124  return (unsigned int) __shfl_up((int) val, delta, width);
125  #else
126  return 0;
127  #endif
128  }
129  __device__ __forceinline__ double shfl_up(double val, unsigned int delta, int width = warpSize)
130  {
131  #if __CUDA_ARCH__ >= 300
132  int lo = __double2loint(val);
133  int hi = __double2hiint(val);
134 
135  lo = __shfl_up(lo, delta, width);
136  hi = __shfl_up(hi, delta, width);
137 
138  return __hiloint2double(hi, lo);
139  #else
140  return 0.0;
141  #endif
142  }
143 }}}
144 
145 #endif // __OPENCV_GPU_WARP_SHUFFLE_HPP__
__device__ __forceinline__ T shfl(T val, int srcLane, int width=warpSize)
Definition: warp_shuffle.hpp:49
GLenum GLsizei width
__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
GLuint GLfloat * val
__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