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_CUDA_WARP_SHUFFLE_HPP__ 44 #define __OPENCV_CUDA_WARP_SHUFFLE_HPP__ 45 46 /** @file 47 * @deprecated Use @ref cudev instead. 48 */ 49 50 //! @cond IGNORED 51 52 namespace cv { namespace cuda { namespace device 53 { 54 template <typename T> shfl(T val,int srcLane,int width=warpSize)55 __device__ __forceinline__ T shfl(T val, int srcLane, int width = warpSize) 56 { 57 #if __CUDA_ARCH__ >= 300 58 return __shfl(val, srcLane, width); 59 #else 60 return T(); 61 #endif 62 } shfl(unsigned int val,int srcLane,int width=warpSize)63 __device__ __forceinline__ unsigned int shfl(unsigned int val, int srcLane, int width = warpSize) 64 { 65 #if __CUDA_ARCH__ >= 300 66 return (unsigned int) __shfl((int) val, srcLane, width); 67 #else 68 return 0; 69 #endif 70 } shfl(double val,int srcLane,int width=warpSize)71 __device__ __forceinline__ double shfl(double val, int srcLane, int width = warpSize) 72 { 73 #if __CUDA_ARCH__ >= 300 74 int lo = __double2loint(val); 75 int hi = __double2hiint(val); 76 77 lo = __shfl(lo, srcLane, width); 78 hi = __shfl(hi, srcLane, width); 79 80 return __hiloint2double(hi, lo); 81 #else 82 return 0.0; 83 #endif 84 } 85 86 template <typename T> shfl_down(T val,unsigned int delta,int width=warpSize)87 __device__ __forceinline__ T shfl_down(T val, unsigned int delta, int width = warpSize) 88 { 89 #if __CUDA_ARCH__ >= 300 90 return __shfl_down(val, delta, width); 91 #else 92 return T(); 93 #endif 94 } shfl_down(unsigned int val,unsigned int delta,int width=warpSize)95 __device__ __forceinline__ unsigned int shfl_down(unsigned int val, unsigned int delta, int width = warpSize) 96 { 97 #if __CUDA_ARCH__ >= 300 98 return (unsigned int) __shfl_down((int) val, delta, width); 99 #else 100 return 0; 101 #endif 102 } shfl_down(double val,unsigned int delta,int width=warpSize)103 __device__ __forceinline__ double shfl_down(double val, unsigned int delta, int width = warpSize) 104 { 105 #if __CUDA_ARCH__ >= 300 106 int lo = __double2loint(val); 107 int hi = __double2hiint(val); 108 109 lo = __shfl_down(lo, delta, width); 110 hi = __shfl_down(hi, delta, width); 111 112 return __hiloint2double(hi, lo); 113 #else 114 return 0.0; 115 #endif 116 } 117 118 template <typename T> shfl_up(T val,unsigned int delta,int width=warpSize)119 __device__ __forceinline__ T shfl_up(T val, unsigned int delta, int width = warpSize) 120 { 121 #if __CUDA_ARCH__ >= 300 122 return __shfl_up(val, delta, width); 123 #else 124 return T(); 125 #endif 126 } shfl_up(unsigned int val,unsigned int delta,int width=warpSize)127 __device__ __forceinline__ unsigned int shfl_up(unsigned int val, unsigned int delta, int width = warpSize) 128 { 129 #if __CUDA_ARCH__ >= 300 130 return (unsigned int) __shfl_up((int) val, delta, width); 131 #else 132 return 0; 133 #endif 134 } shfl_up(double val,unsigned int delta,int width=warpSize)135 __device__ __forceinline__ double shfl_up(double val, unsigned int delta, int width = warpSize) 136 { 137 #if __CUDA_ARCH__ >= 300 138 int lo = __double2loint(val); 139 int hi = __double2hiint(val); 140 141 lo = __shfl_up(lo, delta, width); 142 hi = __shfl_up(hi, delta, width); 143 144 return __hiloint2double(hi, lo); 145 #else 146 return 0.0; 147 #endif 148 } 149 }}} 150 151 //! @endcond 152 153 #endif // __OPENCV_CUDA_WARP_SHUFFLE_HPP__ 154