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 "../lut.hpp" 52 53 #include "opencv2/cudaarithm.hpp" 54 #include "opencv2/cudev.hpp" 55 #include "opencv2/core/private.cuda.hpp" 56 57 using namespace cv; 58 using namespace cv::cuda; 59 using namespace cv::cudev; 60 61 namespace cv { namespace cuda { 62 63 texture<uchar, cudaTextureType1D, cudaReadModeElementType> texLutTable; 64 LookUpTableImpl(InputArray _lut)65 LookUpTableImpl::LookUpTableImpl(InputArray _lut) 66 { 67 if (_lut.kind() == _InputArray::CUDA_GPU_MAT) 68 { 69 d_lut = _lut.getGpuMat(); 70 } 71 else 72 { 73 Mat h_lut = _lut.getMat(); 74 d_lut.upload(Mat(1, 256, h_lut.type(), h_lut.data)); 75 } 76 77 CV_Assert( d_lut.depth() == CV_8U ); 78 CV_Assert( d_lut.rows == 1 && d_lut.cols == 256 ); 79 80 cc30 = deviceSupports(FEATURE_SET_COMPUTE_30); 81 82 if (cc30) 83 { 84 // Use the texture object 85 cudaResourceDesc texRes; 86 std::memset(&texRes, 0, sizeof(texRes)); 87 texRes.resType = cudaResourceTypeLinear; 88 texRes.res.linear.devPtr = d_lut.data; 89 texRes.res.linear.desc = cudaCreateChannelDesc<uchar>(); 90 texRes.res.linear.sizeInBytes = 256 * d_lut.channels() * sizeof(uchar); 91 92 cudaTextureDesc texDescr; 93 std::memset(&texDescr, 0, sizeof(texDescr)); 94 95 CV_CUDEV_SAFE_CALL( cudaCreateTextureObject(&texLutTableObj, &texRes, &texDescr, 0) ); 96 } 97 else 98 { 99 // Use the texture reference 100 cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar>(); 101 CV_CUDEV_SAFE_CALL( cudaBindTexture(0, &texLutTable, d_lut.data, &desc) ); 102 } 103 } 104 ~LookUpTableImpl()105 LookUpTableImpl::~LookUpTableImpl() 106 { 107 if (cc30) 108 { 109 // Use the texture object 110 cudaDestroyTextureObject(texLutTableObj); 111 } 112 else 113 { 114 // Use the texture reference 115 cudaUnbindTexture(texLutTable); 116 } 117 } 118 119 struct LutTablePtrC1 120 { 121 typedef uchar value_type; 122 typedef uchar index_type; 123 124 cudaTextureObject_t texLutTableObj; 125 operator ()cv::cuda::LutTablePtrC1126 __device__ __forceinline__ uchar operator ()(uchar, uchar x) const 127 { 128 #if CV_CUDEV_ARCH < 300 129 // Use the texture reference 130 return tex1Dfetch(texLutTable, x); 131 #else 132 // Use the texture object 133 return tex1Dfetch<uchar>(texLutTableObj, x); 134 #endif 135 } 136 }; 137 struct LutTablePtrC3 138 { 139 typedef uchar3 value_type; 140 typedef uchar3 index_type; 141 142 cudaTextureObject_t texLutTableObj; 143 operator ()cv::cuda::LutTablePtrC3144 __device__ __forceinline__ uchar3 operator ()(const uchar3&, const uchar3& x) const 145 { 146 #if CV_CUDEV_ARCH < 300 147 // Use the texture reference 148 return make_uchar3(tex1Dfetch(texLutTable, x.x * 3), tex1Dfetch(texLutTable, x.y * 3 + 1), tex1Dfetch(texLutTable, x.z * 3 + 2)); 149 #else 150 // Use the texture object 151 return make_uchar3(tex1Dfetch<uchar>(texLutTableObj, x.x * 3), tex1Dfetch<uchar>(texLutTableObj, x.y * 3 + 1), tex1Dfetch<uchar>(texLutTableObj, x.z * 3 + 2)); 152 #endif 153 } 154 }; 155 transform(InputArray _src,OutputArray _dst,Stream & stream)156 void LookUpTableImpl::transform(InputArray _src, OutputArray _dst, Stream& stream) 157 { 158 GpuMat src = getInputMat(_src, stream); 159 160 const int cn = src.channels(); 161 const int lut_cn = d_lut.channels(); 162 163 CV_Assert( src.type() == CV_8UC1 || src.type() == CV_8UC3 ); 164 CV_Assert( lut_cn == 1 || lut_cn == cn ); 165 166 GpuMat dst = getOutputMat(_dst, src.size(), src.type(), stream); 167 168 if (lut_cn == 1) 169 { 170 GpuMat_<uchar> src1(src.reshape(1)); 171 GpuMat_<uchar> dst1(dst.reshape(1)); 172 173 LutTablePtrC1 tbl; 174 tbl.texLutTableObj = texLutTableObj; 175 176 dst1.assign(lut_(src1, tbl), stream); 177 } 178 else if (lut_cn == 3) 179 { 180 GpuMat_<uchar3>& src3 = (GpuMat_<uchar3>&) src; 181 GpuMat_<uchar3>& dst3 = (GpuMat_<uchar3>&) dst; 182 183 LutTablePtrC3 tbl; 184 tbl.texLutTableObj = texLutTableObj; 185 186 dst3.assign(lut_(src3, tbl), stream); 187 } 188 189 syncOutput(dst, _dst, stream); 190 } 191 192 } } 193 194 #endif 195