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 #include "opencv2/core/private.cuda.hpp" 53 54 using namespace cv::cudev; 55 56 void bitScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int op); 57 58 namespace 59 { 60 template <template <typename> class Op, typename T> 61 void bitScalarOp(const GpuMat& src, uint value, GpuMat& dst, Stream& stream) 62 { 63 gridTransformUnary(globPtr<T>(src), globPtr<T>(dst), bind2nd(Op<T>(), value), stream); 64 } 65 66 typedef void (*bit_scalar_func_t)(const GpuMat& src, uint value, GpuMat& dst, Stream& stream); 67 68 template <typename T, bit_scalar_func_t func> struct BitScalar 69 { 70 static void call(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& stream) 71 { 72 func(src, cv::saturate_cast<T>(value[0]), dst, stream); 73 } 74 }; 75 76 template <bit_scalar_func_t func> struct BitScalar4 77 { 78 static void call(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& stream) 79 { 80 uint packedVal = 0; 81 82 packedVal |= cv::saturate_cast<uchar>(value[0]); 83 packedVal |= cv::saturate_cast<uchar>(value[1]) << 8; 84 packedVal |= cv::saturate_cast<uchar>(value[2]) << 16; 85 packedVal |= cv::saturate_cast<uchar>(value[3]) << 24; 86 87 func(src, packedVal, dst, stream); 88 } 89 }; 90 91 template <int DEPTH, int cn> struct NppBitwiseCFunc 92 { 93 typedef typename NPPTypeTraits<DEPTH>::npp_type npp_type; 94 95 typedef NppStatus (*func_t)(const npp_type* pSrc1, int nSrc1Step, const npp_type* pConstants, npp_type* pDst, int nDstStep, NppiSize oSizeROI); 96 }; 97 98 template <int DEPTH, int cn, typename NppBitwiseCFunc<DEPTH, cn>::func_t func> struct NppBitwiseC 99 { 100 typedef typename NppBitwiseCFunc<DEPTH, cn>::npp_type npp_type; 101 102 static void call(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& _stream) 103 { 104 cudaStream_t stream = StreamAccessor::getStream(_stream); 105 NppStreamHandler h(stream); 106 107 NppiSize oSizeROI; 108 oSizeROI.width = src.cols; 109 oSizeROI.height = src.rows; 110 111 const npp_type pConstants[] = 112 { 113 cv::saturate_cast<npp_type>(value[0]), 114 cv::saturate_cast<npp_type>(value[1]), 115 cv::saturate_cast<npp_type>(value[2]), 116 cv::saturate_cast<npp_type>(value[3]) 117 }; 118 119 nppSafeCall( func(src.ptr<npp_type>(), static_cast<int>(src.step), pConstants, dst.ptr<npp_type>(), static_cast<int>(dst.step), oSizeROI) ); 120 121 if (stream == 0) 122 CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); 123 } 124 }; 125 } 126 127 void bitScalar(const GpuMat& src, cv::Scalar value, bool, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int op) 128 { 129 (void) mask; 130 131 typedef void (*func_t)(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& stream); 132 static const func_t funcs[3][6][4] = 133 { 134 { 135 {BitScalar<uchar, bitScalarOp<bit_and, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiAndC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_and, uint> >::call}, 136 {BitScalar<uchar, bitScalarOp<bit_and, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiAndC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_and, uint> >::call}, 137 {BitScalar<ushort, bitScalarOp<bit_and, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiAndC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiAndC_16u_C4R>::call}, 138 {BitScalar<ushort, bitScalarOp<bit_and, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiAndC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiAndC_16u_C4R>::call}, 139 {BitScalar<uint, bitScalarOp<bit_and, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiAndC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiAndC_32s_C4R>::call}, 140 {BitScalar<uint, bitScalarOp<bit_and, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiAndC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiAndC_32s_C4R>::call} 141 }, 142 { 143 {BitScalar<uchar, bitScalarOp<bit_or, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiOrC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_or, uint> >::call}, 144 {BitScalar<uchar, bitScalarOp<bit_or, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiOrC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_or, uint> >::call}, 145 {BitScalar<ushort, bitScalarOp<bit_or, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiOrC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiOrC_16u_C4R>::call}, 146 {BitScalar<ushort, bitScalarOp<bit_or, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiOrC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiOrC_16u_C4R>::call}, 147 {BitScalar<uint, bitScalarOp<bit_or, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiOrC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiOrC_32s_C4R>::call}, 148 {BitScalar<uint, bitScalarOp<bit_or, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiOrC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiOrC_32s_C4R>::call} 149 }, 150 { 151 {BitScalar<uchar, bitScalarOp<bit_xor, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiXorC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_xor, uint> >::call}, 152 {BitScalar<uchar, bitScalarOp<bit_xor, uchar> >::call , 0, NppBitwiseC<CV_8U , 3, nppiXorC_8u_C3R >::call, BitScalar4< bitScalarOp<bit_xor, uint> >::call}, 153 {BitScalar<ushort, bitScalarOp<bit_xor, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiXorC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiXorC_16u_C4R>::call}, 154 {BitScalar<ushort, bitScalarOp<bit_xor, ushort> >::call, 0, NppBitwiseC<CV_16U, 3, nppiXorC_16u_C3R>::call, NppBitwiseC<CV_16U, 4, nppiXorC_16u_C4R>::call}, 155 {BitScalar<uint, bitScalarOp<bit_xor, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiXorC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiXorC_32s_C4R>::call}, 156 {BitScalar<uint, bitScalarOp<bit_xor, uint> >::call , 0, NppBitwiseC<CV_32S, 3, nppiXorC_32s_C3R>::call, NppBitwiseC<CV_32S, 4, nppiXorC_32s_C4R>::call} 157 } 158 }; 159 160 const int depth = src.depth(); 161 const int cn = src.channels(); 162 163 CV_DbgAssert( depth <= CV_32F ); 164 CV_DbgAssert( cn == 1 || cn == 3 || cn == 4 ); 165 CV_DbgAssert( mask.empty() ); 166 CV_DbgAssert( op >= 0 && op < 3 ); 167 168 funcs[op][depth][cn - 1](src, value, dst, stream); 169 } 170 171 #endif 172
