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>
bitScalarOp(const GpuMat & src,uint value,GpuMat & dst,Stream & stream)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     {
call__anon67c1016c0111::BitScalar70         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     {
call__anon67c1016c0111::BitScalar478         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 
call__anon67c1016c0111::NppBitwiseC102         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 
bitScalar(const GpuMat & src,cv::Scalar value,bool,GpuMat & dst,const GpuMat & mask,double,Stream & stream,int op)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