From 7316676c41033ba62daa6780c8ef1307ab568205 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 28 Aug 2014 14:47:26 +0400 Subject: [PATCH] fix CUDA LUT implementation In CUDA 6.0 there was a bug in NPP LUT implementation (invalid results when src == 255). In CUDA 6.5 the bug was fixed. Replaced NPP LUT call with own implementation (ported from master branch) to be independant from CUDA Toolkit version. (cherry picked from commit eaaa2d27d5ab334c74c2d10550a6097f437fb297) --- modules/gpu/src/arithm.cpp | 82 ++++---------------- modules/gpu/src/cuda/lut.cu | 151 ++++++++++++++++++++++++++++++++++++ 2 files changed, 164 insertions(+), 69 deletions(-) create mode 100644 modules/gpu/src/cuda/lut.cu diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 53b6aea0c5..df001d037b 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -317,6 +317,11 @@ void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode, Stream& stream) //////////////////////////////////////////////////////////////////////// // LUT +namespace arithm +{ + void lut(PtrStepSzb src, uchar* lut, int lut_cn, PtrStepSzb dst, bool cc30, cudaStream_t stream); +} + void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) { const int cn = src.channels(); @@ -328,82 +333,21 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) dst.create(src.size(), CV_MAKE_TYPE(lut.depth(), cn)); - NppiSize sz; - sz.height = src.rows; - sz.width = src.cols; - - Mat nppLut; - lut.convertTo(nppLut, CV_32S); - - int nValues3[] = {256, 256, 256}; - - Npp32s pLevels[256]; - for (int i = 0; i < 256; ++i) - pLevels[i] = i; - - const Npp32s* pLevels3[3]; - -#if (CUDA_VERSION <= 4020) - pLevels3[0] = pLevels3[1] = pLevels3[2] = pLevels; -#else - GpuMat d_pLevels; - d_pLevels.upload(Mat(1, 256, CV_32S, pLevels)); - pLevels3[0] = pLevels3[1] = pLevels3[2] = d_pLevels.ptr(); -#endif + GpuMat d_lut; + d_lut.upload(Mat(1, 256, lut.type(), lut.data)); + int lut_cn = d_lut.channels(); + bool cc30 = deviceSupports(FEATURE_SET_COMPUTE_30); cudaStream_t stream = StreamAccessor::getStream(s); - NppStreamHandler h(stream); - if (src.type() == CV_8UC1) + if (lut_cn == 1) { -#if (CUDA_VERSION <= 4020) - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, nppLut.ptr(), pLevels, 256) ); -#else - GpuMat d_nppLut(Mat(1, 256, CV_32S, nppLut.data)); - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, d_nppLut.ptr(), d_pLevels.ptr(), 256) ); -#endif + arithm::lut(src.reshape(1), d_lut.data, lut_cn, dst.reshape(1), cc30, stream); } - else + else if (lut_cn == 3) { - const Npp32s* pValues3[3]; - - Mat nppLut3[3]; - if (nppLut.channels() == 1) - { -#if (CUDA_VERSION <= 4020) - pValues3[0] = pValues3[1] = pValues3[2] = nppLut.ptr(); -#else - GpuMat d_nppLut(Mat(1, 256, CV_32S, nppLut.data)); - pValues3[0] = pValues3[1] = pValues3[2] = d_nppLut.ptr(); -#endif - } - else - { - cv::split(nppLut, nppLut3); - -#if (CUDA_VERSION <= 4020) - pValues3[0] = nppLut3[0].ptr(); - pValues3[1] = nppLut3[1].ptr(); - pValues3[2] = nppLut3[2].ptr(); -#else - GpuMat d_nppLut0(Mat(1, 256, CV_32S, nppLut3[0].data)); - GpuMat d_nppLut1(Mat(1, 256, CV_32S, nppLut3[1].data)); - GpuMat d_nppLut2(Mat(1, 256, CV_32S, nppLut3[2].data)); - - pValues3[0] = d_nppLut0.ptr(); - pValues3[1] = d_nppLut1.ptr(); - pValues3[2] = d_nppLut2.ptr(); -#endif - } - - nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, pValues3, pLevels3, nValues3) ); + arithm::lut(src, d_lut.data, lut_cn, dst, cc30, stream); } - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); } //////////////////////////////////////////////////////////////////////// diff --git a/modules/gpu/src/cuda/lut.cu b/modules/gpu/src/cuda/lut.cu new file mode 100644 index 0000000000..be5efeca2b --- /dev/null +++ b/modules/gpu/src/cuda/lut.cu @@ -0,0 +1,151 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#if !defined CUDA_DISABLER + +#include +#include "opencv2/gpu/device/common.hpp" +#include "opencv2/gpu/device/transform.hpp" +#include "opencv2/gpu/device/functional.hpp" + +using namespace cv::gpu; +using namespace cv::gpu::device; + +namespace +{ + texture texLutTable; + + struct LutC1 : public unary_function + { + typedef uchar value_type; + typedef uchar index_type; + + cudaTextureObject_t texLutTableObj; + + __device__ __forceinline__ uchar operator ()(uchar x) const + { + #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 300) + // Use the texture reference + return tex1Dfetch(texLutTable, x); + #else + // Use the texture object + return tex1Dfetch(texLutTableObj, x); + #endif + } + }; + struct LutC3 : public unary_function + { + typedef uchar3 value_type; + typedef uchar3 index_type; + + cudaTextureObject_t texLutTableObj; + + __device__ __forceinline__ uchar3 operator ()(const uchar3& x) const + { + #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ < 300) + // Use the texture reference + return make_uchar3(tex1Dfetch(texLutTable, x.x * 3), tex1Dfetch(texLutTable, x.y * 3 + 1), tex1Dfetch(texLutTable, x.z * 3 + 2)); + #else + // Use the texture object + return make_uchar3(tex1Dfetch(texLutTableObj, x.x * 3), tex1Dfetch(texLutTableObj, x.y * 3 + 1), tex1Dfetch(texLutTableObj, x.z * 3 + 2)); + #endif + } + }; +} + +namespace arithm +{ + void lut(PtrStepSzb src, uchar* lut, int lut_cn, PtrStepSzb dst, bool cc30, cudaStream_t stream) + { + cudaTextureObject_t texLutTableObj; + + if (cc30) + { + // Use the texture object + cudaResourceDesc texRes; + std::memset(&texRes, 0, sizeof(texRes)); + texRes.resType = cudaResourceTypeLinear; + texRes.res.linear.devPtr = lut; + texRes.res.linear.desc = cudaCreateChannelDesc(); + texRes.res.linear.sizeInBytes = 256 * lut_cn * sizeof(uchar); + + cudaTextureDesc texDescr; + std::memset(&texDescr, 0, sizeof(texDescr)); + + cudaSafeCall( cudaCreateTextureObject(&texLutTableObj, &texRes, &texDescr, 0) ); + } + else + { + // Use the texture reference + cudaChannelFormatDesc desc = cudaCreateChannelDesc(); + cudaSafeCall( cudaBindTexture(0, &texLutTable, lut, &desc) ); + } + + if (lut_cn == 1) + { + LutC1 op; + op.texLutTableObj = texLutTableObj; + + transform((PtrStepSz) src, (PtrStepSz) dst, op, WithOutMask(), stream); + } + else if (lut_cn == 3) + { + LutC3 op; + op.texLutTableObj = texLutTableObj; + + transform((PtrStepSz) src, (PtrStepSz) dst, op, WithOutMask(), stream); + } + + if (cc30) + { + // Use the texture object + cudaSafeCall( cudaDestroyTextureObject(texLutTableObj) ); + } + else + { + // Use the texture reference + cudaSafeCall( cudaUnbindTexture(texLutTable) ); + } + } +} + +#endif