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 eaaa2d27d5)
This commit is contained in:
Vladislav Vinogradov
Alexander Smorkalov
parent
77585bf8af
commit
7316676c41
@ -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<Npp32s>();
|
||||
#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<Npp8u>(), static_cast<int>(src.step),
|
||||
dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, nppLut.ptr<Npp32s>(), pLevels, 256) );
|
||||
#else
|
||||
GpuMat d_nppLut(Mat(1, 256, CV_32S, nppLut.data));
|
||||
nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step),
|
||||
dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, d_nppLut.ptr<Npp32s>(), d_pLevels.ptr<Npp32s>(), 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<Npp32s>();
|
||||
#else
|
||||
GpuMat d_nppLut(Mat(1, 256, CV_32S, nppLut.data));
|
||||
pValues3[0] = pValues3[1] = pValues3[2] = d_nppLut.ptr<Npp32s>();
|
||||
#endif
|
||||
}
|
||||
else
|
||||
{
|
||||
cv::split(nppLut, nppLut3);
|
||||
|
||||
#if (CUDA_VERSION <= 4020)
|
||||
pValues3[0] = nppLut3[0].ptr<Npp32s>();
|
||||
pValues3[1] = nppLut3[1].ptr<Npp32s>();
|
||||
pValues3[2] = nppLut3[2].ptr<Npp32s>();
|
||||
#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<Npp32s>();
|
||||
pValues3[1] = d_nppLut1.ptr<Npp32s>();
|
||||
pValues3[2] = d_nppLut2.ptr<Npp32s>();
|
||||
#endif
|
||||
}
|
||||
|
||||
nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr<Npp8u>(), static_cast<int>(src.step),
|
||||
dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, pValues3, pLevels3, nValues3) );
|
||||
arithm::lut(src, d_lut.data, lut_cn, dst, cc30, stream);
|
||||
}
|
||||
|
||||
if (stream == 0)
|
||||
cudaSafeCall( cudaDeviceSynchronize() );
|
||||
}
|
||||
|
||||
////////////////////////////////////////////////////////////////////////
|
||||
|
||||
151
modules/gpu/src/cuda/lut.cu
Normal file
151
modules/gpu/src/cuda/lut.cu
Normal file
@ -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 <cstring>
|
||||
#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<uchar, cudaTextureType1D, cudaReadModeElementType> texLutTable;
|
||||
|
||||
struct LutC1 : public unary_function<uchar, uchar>
|
||||
{
|
||||
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<uchar>(texLutTableObj, x);
|
||||
#endif
|
||||
}
|
||||
};
|
||||
struct LutC3 : public unary_function<uchar3, uchar3>
|
||||
{
|
||||
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<uchar>(texLutTableObj, x.x * 3), tex1Dfetch<uchar>(texLutTableObj, x.y * 3 + 1), tex1Dfetch<uchar>(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<uchar>();
|
||||
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<uchar>();
|
||||
cudaSafeCall( cudaBindTexture(0, &texLutTable, lut, &desc) );
|
||||
}
|
||||
|
||||
if (lut_cn == 1)
|
||||
{
|
||||
LutC1 op;
|
||||
op.texLutTableObj = texLutTableObj;
|
||||
|
||||
transform((PtrStepSz<uchar>) src, (PtrStepSz<uchar>) dst, op, WithOutMask(), stream);
|
||||
}
|
||||
else if (lut_cn == 3)
|
||||
{
|
||||
LutC3 op;
|
||||
op.texLutTableObj = texLutTableObj;
|
||||
|
||||
transform((PtrStepSz<uchar3>) src, (PtrStepSz<uchar3>) dst, op, WithOutMask(), stream);
|
||||
}
|
||||
|
||||
if (cc30)
|
||||
{
|
||||
// Use the texture object
|
||||
cudaSafeCall( cudaDestroyTextureObject(texLutTableObj) );
|
||||
}
|
||||
else
|
||||
{
|
||||
// Use the texture reference
|
||||
cudaSafeCall( cudaUnbindTexture(texLutTable) );
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
Loading…
Reference in New Issue
Block a user