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 #include "opencv2/opencv_modules.hpp" 44 45 #ifndef HAVE_OPENCV_CUDEV 46 47 #error "opencv_cudev is required" 48 49 #else 50 51 #include "opencv2/cudev.hpp" 52 53 using namespace cv::cudev; 54 55 void minMaxMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int op); 56 57 void minMaxScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int op); 58 59 /////////////////////////////////////////////////////////////////////// 60 /// minMaxMat 61 62 namespace 63 { 64 template <template <typename> class Op, typename T> 65 void minMaxMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) 66 { 67 gridTransformBinary(globPtr<T>(src1), globPtr<T>(src2), globPtr<T>(dst), Op<T>(), stream); 68 } 69 70 struct MinOp2 : binary_function<uint, uint, uint> 71 { 72 __device__ __forceinline__ uint operator ()(uint a, uint b) const 73 { 74 return vmin2(a, b); 75 } 76 }; 77 78 struct MaxOp2 : binary_function<uint, uint, uint> 79 { 80 __device__ __forceinline__ uint operator ()(uint a, uint b) const 81 { 82 return vmax2(a, b); 83 } 84 }; 85 86 template <class Op2> 87 void minMaxMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) 88 { 89 const int vcols = src1.cols >> 1; 90 91 GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); 92 GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); 93 GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); 94 95 gridTransformBinary(src1_, src2_, dst_, Op2(), stream); 96 } 97 98 struct MinOp4 : binary_function<uint, uint, uint> 99 { 100 __device__ __forceinline__ uint operator ()(uint a, uint b) const 101 { 102 return vmin4(a, b); 103 } 104 }; 105 106 struct MaxOp4 : binary_function<uint, uint, uint> 107 { 108 __device__ __forceinline__ uint operator ()(uint a, uint b) const 109 { 110 return vmax4(a, b); 111 } 112 }; 113 114 template <class Op4> 115 void minMaxMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) 116 { 117 const int vcols = src1.cols >> 2; 118 119 GlobPtrSz<uint> src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); 120 GlobPtrSz<uint> src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); 121 GlobPtrSz<uint> dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); 122 123 gridTransformBinary(src1_, src2_, dst_, Op4(), stream); 124 } 125 } 126 127 void minMaxMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int op) 128 { 129 typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); 130 static const func_t funcs_v1[2][7] = 131 { 132 { 133 minMaxMat_v1<minimum, uchar>, 134 minMaxMat_v1<minimum, schar>, 135 minMaxMat_v1<minimum, ushort>, 136 minMaxMat_v1<minimum, short>, 137 minMaxMat_v1<minimum, int>, 138 minMaxMat_v1<minimum, float>, 139 minMaxMat_v1<minimum, double> 140 }, 141 { 142 minMaxMat_v1<maximum, uchar>, 143 minMaxMat_v1<maximum, schar>, 144 minMaxMat_v1<maximum, ushort>, 145 minMaxMat_v1<maximum, short>, 146 minMaxMat_v1<maximum, int>, 147 minMaxMat_v1<maximum, float>, 148 minMaxMat_v1<maximum, double> 149 } 150 }; 151 152 static const func_t funcs_v2[2] = 153 { 154 minMaxMat_v2<MinOp2>, minMaxMat_v2<MaxOp2> 155 }; 156 157 static const func_t funcs_v4[2] = 158 { 159 minMaxMat_v4<MinOp4>, minMaxMat_v4<MaxOp4> 160 }; 161 162 const int depth = src1.depth(); 163 164 CV_DbgAssert( depth <= CV_64F ); 165 166 GpuMat src1_ = src1.reshape(1); 167 GpuMat src2_ = src2.reshape(1); 168 GpuMat dst_ = dst.reshape(1); 169 170 if (depth == CV_8U || depth == CV_16U) 171 { 172 const intptr_t src1ptr = reinterpret_cast<intptr_t>(src1_.data); 173 const intptr_t src2ptr = reinterpret_cast<intptr_t>(src2_.data); 174 const intptr_t dstptr = reinterpret_cast<intptr_t>(dst_.data); 175 176 const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; 177 178 if (isAllAligned) 179 { 180 if (depth == CV_8U && (src1_.cols & 3) == 0) 181 { 182 funcs_v4[op](src1_, src2_, dst_, stream); 183 return; 184 } 185 else if (depth == CV_16U && (src1_.cols & 1) == 0) 186 { 187 funcs_v2[op](src1_, src2_, dst_, stream); 188 return; 189 } 190 } 191 } 192 193 const func_t func = funcs_v1[op][depth]; 194 195 func(src1_, src2_, dst_, stream); 196 } 197 198 /////////////////////////////////////////////////////////////////////// 199 /// minMaxScalar 200 201 namespace 202 { 203 template <template <typename> class Op, typename T> 204 void minMaxScalar(const GpuMat& src, double value, GpuMat& dst, Stream& stream) 205 { 206 gridTransformUnary(globPtr<T>(src), globPtr<T>(dst), bind2nd(Op<T>(), cv::saturate_cast<T>(value)), stream); 207 } 208 } 209 210 void minMaxScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int op) 211 { 212 typedef void (*func_t)(const GpuMat& src, double value, GpuMat& dst, Stream& stream); 213 static const func_t funcs[2][7] = 214 { 215 { 216 minMaxScalar<minimum, uchar>, 217 minMaxScalar<minimum, schar>, 218 minMaxScalar<minimum, ushort>, 219 minMaxScalar<minimum, short>, 220 minMaxScalar<minimum, int>, 221 minMaxScalar<minimum, float>, 222 minMaxScalar<minimum, double> 223 }, 224 { 225 minMaxScalar<maximum, uchar>, 226 minMaxScalar<maximum, schar>, 227 minMaxScalar<maximum, ushort>, 228 minMaxScalar<maximum, short>, 229 minMaxScalar<maximum, int>, 230 minMaxScalar<maximum, float>, 231 minMaxScalar<maximum, double> 232 } 233 }; 234 235 const int depth = src.depth(); 236 237 CV_DbgAssert( depth <= CV_64F ); 238 CV_DbgAssert( src.channels() == 1 ); 239 240 funcs[op][depth](src, value[0], dst, stream); 241 } 242 243 #endif 244
