Skip to content

O
opencv
Commit 80c1aecf authored by Vladislav Vinogradov's avatar Vladislav Vinogradov
Browse files
Options

moved gpu::cvtColor from imgproc_gpu.cpp to color.cpp

parent 949d3109
Show whitespace changes
Inline Side-by-side
Showing with 524 additions and 474 deletions
modules/gpu/src/color.cpp 0 → 100644
View file @ 80c1aecf
/*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*/
#include "precomp.hpp"
using namespace cv;
using namespace cv::gpu;
#if !defined (HAVE_CUDA)
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }
#else /* !defined (HAVE_CUDA) */
namespace cv { namespace gpu { namespace color
{
void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream);
void Gray2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream);
void RGB2Gray_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB2Gray_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream);
void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HSV_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HSV2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HLS_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HLS2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
}}}
namespace
{
#undef R2Y
#undef G2Y
#undef B2Y
enum
{
yuv_shift = 14,
xyz_shift = 12,
R2Y = 4899,
G2Y = 9617,
B2Y = 1868,
BLOCK_SIZE = 256
};
}
namespace
{
void cvtColor_caller(const GpuMat& src, GpuMat& dst, int code, int dcn, const cudaStream_t& stream)
{
Size sz = src.size();
int scn = src.channels(), depth = src.depth(), bidx;
CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F);
switch (code)
{
case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
static const func_t funcs[] = {color::RGB2RGB_gpu_8u, 0, color::RGB2RGB_gpu_16u, 0, 0, color::RGB2RGB_gpu_32f};
CV_Assert(scn == 3 || scn == 4);
dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3;
bidx = code == CV_BGR2BGRA || code == CV_BGRA2BGR ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, dcn));
funcs[depth](src, scn, dst, dcn, bidx, stream);
break;
}
case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
{
CV_Assert((scn == 3 || scn == 4) && depth == CV_8U);
int green_bits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565
|| code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;
bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555
|| code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2;
dst.create(sz, CV_8UC2);
color::RGB2RGB5x5_gpu(src, scn, dst, green_bits, bidx, stream);
break;
}
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
{
if (dcn <= 0) dcn = 3;
CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);
int green_bits = code == CV_BGR5652BGR || code == CV_BGR5652RGB
|| code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;
bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR
|| code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, dcn));
color::RGB5x52RGB_gpu(src, green_bits, dst, dcn, bidx, stream);
break;
}
case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
static const func_t funcs[] = {color::RGB2Gray_gpu_8u, 0, color::RGB2Gray_gpu_16u, 0, 0, color::RGB2Gray_gpu_32f};
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 1));
funcs[depth](src, scn, dst, bidx, stream);
break;
}
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
{
CV_Assert(scn == 2 && depth == CV_8U);
int green_bits = code == CV_BGR5652GRAY ? 6 : 5;
dst.create(sz, CV_8UC1);
color::RGB5x52Gray_gpu(src, green_bits, dst, stream);
break;
}
case CV_GRAY2BGR: case CV_GRAY2BGRA:
{
typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
static const func_t funcs[] = {color::Gray2RGB_gpu_8u, 0, color::Gray2RGB_gpu_16u, 0, 0, color::Gray2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert(scn == 1 && (dcn == 3 || dcn == 4));
dst.create(sz, CV_MAKETYPE(depth, dcn));
funcs[depth](src, dst, dcn, stream);
break;
}
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
{
CV_Assert(scn == 1 && depth == CV_8U);
int green_bits = code == CV_GRAY2BGR565 ? 6 : 5;
dst.create(sz, CV_8UC2);
color::Gray2RGB5x5_gpu(src, dst, green_bits, stream);
break;
}
case CV_BGR2YCrCb: case CV_RGB2YCrCb:
case CV_BGR2YUV: case CV_RGB2YUV:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {color::RGB2YCrCb_gpu_8u, 0, color::RGB2YCrCb_gpu_16u, 0, 0, color::RGB2YCrCb_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_BGR2YCrCb || code == CV_RGB2YUV ? 0 : 2;
static const float yuv_f[] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
static const int yuv_i[] = { B2Y, G2Y, R2Y, 8061, 14369 };
static const float YCrCb_f[] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
static const int YCrCb_i[] = {R2Y, G2Y, B2Y, 11682, 9241};
float coeffs_f[5];
int coeffs_i[5];
::memcpy(coeffs_f, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_f : yuv_f, sizeof(yuv_f));
::memcpy(coeffs_i, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_i : yuv_i, sizeof(yuv_i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[2]);
std::swap(coeffs_i[0], coeffs_i[2]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);
break;
}
case CV_YCrCb2BGR: case CV_YCrCb2RGB:
case CV_YUV2BGR: case CV_YUV2RGB:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {color::YCrCb2RGB_gpu_8u, 0, color::YCrCb2RGB_gpu_16u, 0, 0, color::YCrCb2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_YCrCb2BGR || code == CV_YUV2RGB ? 0 : 2;
static const float yuv_f[] = { 2.032f, -0.395f, -0.581f, 1.140f };
static const int yuv_i[] = { 33292, -6472, -9519, 18678 };
static const float YCrCb_f[] = {1.403f, -0.714f, -0.344f, 1.773f};
static const int YCrCb_i[] = {22987, -11698, -5636, 29049};
const float* coeffs_f = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_f : yuv_f;
const int* coeffs_i = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_i : yuv_i;
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);
break;
}
case CV_BGR2XYZ: case CV_RGB2XYZ:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {color::RGB2XYZ_gpu_8u, 0, color::RGB2XYZ_gpu_16u, 0, 0, color::RGB2XYZ_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_BGR2XYZ ? 0 : 2;
static const float RGB2XYZ_D65f[] =
{
0.412453f, 0.357580f, 0.180423f,
0.212671f, 0.715160f, 0.072169f,
0.019334f, 0.119193f, 0.950227f
};
static const int RGB2XYZ_D65i[] =
{
1689, 1465, 739,
871, 2929, 296,
79, 488, 3892
};
float coeffs_f[9];
int coeffs_i[9];
::memcpy(coeffs_f, RGB2XYZ_D65f, sizeof(RGB2XYZ_D65f));
::memcpy(coeffs_i, RGB2XYZ_D65i, sizeof(RGB2XYZ_D65i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[2]);
std::swap(coeffs_f[3], coeffs_f[5]);
std::swap(coeffs_f[6], coeffs_f[8]);
std::swap(coeffs_i[0], coeffs_i[2]);
std::swap(coeffs_i[3], coeffs_i[5]);
std::swap(coeffs_i[6], coeffs_i[8]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, coeffs, stream);
break;
}
case CV_XYZ2BGR: case CV_XYZ2RGB:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {color::XYZ2RGB_gpu_8u, 0, color::XYZ2RGB_gpu_16u, 0, 0, color::XYZ2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_XYZ2BGR ? 0 : 2;
static const float XYZ2sRGB_D65f[] =
{
3.240479f, -1.53715f, -0.498535f,
-0.969256f, 1.875991f, 0.041556f,
0.055648f, -0.204043f, 1.057311f
};
static const int XYZ2sRGB_D65i[] =
{
13273, -6296, -2042,
-3970, 7684, 170,
228, -836, 4331
};
float coeffs_f[9];
int coeffs_i[9];
::memcpy(coeffs_f, XYZ2sRGB_D65f, sizeof(XYZ2sRGB_D65f));
::memcpy(coeffs_i, XYZ2sRGB_D65i, sizeof(XYZ2sRGB_D65i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[6]);
std::swap(coeffs_f[1], coeffs_f[7]);
std::swap(coeffs_f[2], coeffs_f[8]);
std::swap(coeffs_i[0], coeffs_i[6]);
std::swap(coeffs_i[1], coeffs_i[7]);
std::swap(coeffs_i[2], coeffs_i[8]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, coeffs_i, stream);
break;
}
case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:
case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
int hrange, cudaStream_t stream);
static const func_t funcs_hsv[] = {color::RGB2HSV_gpu_8u, 0, 0, 0, 0, color::RGB2HSV_gpu_32f};
static const func_t funcs_hls[] = {color::RGB2HLS_gpu_8u, 0, 0, 0, 0, color::RGB2HLS_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == CV_BGR2HSV || code == CV_BGR2HLS ||
code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV ||
code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 255;
dst.create(sz, CV_MAKETYPE(depth, dcn));
if (code == CV_BGR2HSV || code == CV_RGB2HSV || code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL)
funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);
else
funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);
break;
}
case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:
case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
int hrange, cudaStream_t stream);
static const func_t funcs_hsv[] = {color::HSV2RGB_gpu_8u, 0, 0, 0, 0, color::HSV2RGB_gpu_32f};
static const func_t funcs_hls[] = {color::HLS2RGB_gpu_8u, 0, 0, 0, 0, color::HLS2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == CV_HSV2BGR || code == CV_HLS2BGR ||
code == CV_HSV2BGR_FULL || code == CV_HLS2BGR_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == CV_HSV2BGR || code == CV_HSV2RGB ||
code == CV_HLS2BGR || code == CV_HLS2RGB ? 180 : 255;
dst.create(sz, CV_MAKETYPE(depth, dcn));
if (code == CV_HSV2BGR || code == CV_HSV2RGB || code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL)
funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);
else
funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);
break;
}
default:
CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
}
}
}
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)
{
cvtColor_caller(src, dst, code, dcn, 0);
}
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)
{
cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));
}
#endif /* !defined (HAVE_CUDA) */
modules/gpu/src/cuda/color.cu
View file @ 80c1aecf
......@@ -54,7 +54,7 @@ using namespace cv::gpu;
#define FLT_EPSILON 1.192092896e-07F
#endif
namespace imgproc_krnls
namespace color_krnls
{
template<typename T> struct ColorChannel {};
template<> struct ColorChannel<uchar>
......@@ -99,7 +99,7 @@ namespace imgproc_krnls
////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////
namespace imgproc_krnls
namespace color_krnls
{
template <int SRCCN, int DSTCN, typename T>
__global__ void RGB2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)
......@@ -125,7 +125,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN, int DSTCN>
void RGB2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)
......@@ -136,7 +136,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -182,7 +182,7 @@ namespace cv { namespace gpu { namespace imgproc
/////////// Transforming 16-bit (565 or 555) RGB to/from 24/32-bit (888[8]) RGB //////////
namespace imgproc_krnls
namespace color_krnls
{
template <int GREEN_BITS, int DSTCN> struct RGB5x52RGBConverter {};
template <int DSTCN> struct RGB5x52RGBConverter<5, DSTCN>
......@@ -274,7 +274,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <int GREEN_BITS, int DSTCN>
void RGB5x52RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)
......@@ -285,7 +285,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB5x52RGB<GREEN_BITS, DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB5x52RGB<GREEN_BITS, DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -313,7 +313,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB2RGB5x5<SRCCN, GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2RGB5x5<SRCCN, GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -335,7 +335,7 @@ namespace cv { namespace gpu { namespace imgproc
///////////////////////////////// Grayscale to Color ////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
template <int DSTCN, typename T>
__global__ void Gray2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols)
......@@ -389,7 +389,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int DSTCN>
void Gray2RGB_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)
......@@ -400,7 +400,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::Gray2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::Gray2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols);
if (stream == 0)
......@@ -440,7 +440,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::Gray2RGB5x5<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::Gray2RGB5x5<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols);
if (stream == 0)
......@@ -461,7 +461,7 @@ namespace cv { namespace gpu { namespace imgproc
///////////////////////////////// Color to Grayscale ////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
#undef R2Y
#undef G2Y
......@@ -543,7 +543,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN>
void RGB2Gray_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)
......@@ -554,7 +554,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB2Gray<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2Gray<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -594,7 +594,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB5x52Gray<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB5x52Gray<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols);
if (stream == 0)
......@@ -615,7 +615,7 @@ namespace cv { namespace gpu { namespace imgproc
///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
__constant__ float cYCrCbCoeffs_f[5];
__constant__ int cYCrCbCoeffs_i[5];
......@@ -714,7 +714,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN, int DSTCN>
void RGB2YCrCb_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)
......@@ -725,7 +725,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB2YCrCb<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2YCrCb<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -741,7 +741,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2YCrCb_caller<uchar, 4, 3>, RGB2YCrCb_caller<uchar, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -755,7 +755,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2YCrCb_caller<ushort, 4, 3>, RGB2YCrCb_caller<ushort, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -769,7 +769,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2YCrCb_caller<float, 4, 3>, RGB2YCrCb_caller<float, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) );
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -783,7 +783,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::YCrCb2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::YCrCb2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -799,7 +799,7 @@ namespace cv { namespace gpu { namespace imgproc
{YCrCb2RGB_caller<uchar, 4, 3>, YCrCb2RGB_caller<uchar, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -813,7 +813,7 @@ namespace cv { namespace gpu { namespace imgproc
{YCrCb2RGB_caller<ushort, 4, 3>, YCrCb2RGB_caller<ushort, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -827,7 +827,7 @@ namespace cv { namespace gpu { namespace imgproc
{YCrCb2RGB_caller<float, 4, 3>, YCrCb2RGB_caller<float, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) );
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);
}
......@@ -835,7 +835,7 @@ namespace cv { namespace gpu { namespace imgproc
////////////////////////////////////// RGB <-> XYZ ///////////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
__constant__ float cXYZ_D65f[9];
__constant__ int cXYZ_D65i[9];
......@@ -924,7 +924,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN, int DSTCN>
void RGB2XYZ_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)
......@@ -935,7 +935,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::RGB2XYZ<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2XYZ<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols);
if (stream == 0)
......@@ -951,7 +951,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2XYZ_caller<uchar, 4, 3>, RGB2XYZ_caller<uchar, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -965,7 +965,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2XYZ_caller<ushort, 4, 3>, RGB2XYZ_caller<ushort, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -979,7 +979,7 @@ namespace cv { namespace gpu { namespace imgproc
{RGB2XYZ_caller<float, 4, 3>, RGB2XYZ_caller<float, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -993,7 +993,7 @@ namespace cv { namespace gpu { namespace imgproc
grid.x = divUp(src.cols, threads.x);
grid.y = divUp(src.rows, threads.y);
imgproc_krnls::XYZ2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::XYZ2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols);
if (stream == 0)
......@@ -1009,7 +1009,7 @@ namespace cv { namespace gpu { namespace imgproc
{XYZ2RGB_caller<uchar, 4, 3>, XYZ2RGB_caller<uchar, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -1023,7 +1023,7 @@ namespace cv { namespace gpu { namespace imgproc
{XYZ2RGB_caller<ushort, 4, 3>, XYZ2RGB_caller<ushort, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -1037,7 +1037,7 @@ namespace cv { namespace gpu { namespace imgproc
{XYZ2RGB_caller<float, 4, 3>, XYZ2RGB_caller<float, 4, 4>}
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);
}
......@@ -1045,7 +1045,7 @@ namespace cv { namespace gpu { namespace imgproc
////////////////////////////////////// RGB <-> HSV ///////////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
__constant__ int cHsvDivTable[256];
......@@ -1222,7 +1222,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN, int DSTCN>
void RGB2HSV_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
......@@ -1234,10 +1234,10 @@ namespace cv { namespace gpu { namespace imgproc
grid.y = divUp(src.rows, threads.y);
if (hrange == 180)
imgproc_krnls::RGB2HSV<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2HSV<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
else
imgproc_krnls::RGB2HSV<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2HSV<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -1288,7 +1288,7 @@ namespace cv { namespace gpu { namespace imgproc
4352, 4334, 4316, 4298, 4281, 4263, 4246, 4229,
4212, 4195, 4178, 4161, 4145, 4128, 4112, 4096
};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvDivTable, div_table, sizeof(div_table)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvDivTable, div_table, sizeof(div_table)) );
RGB2HSV_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);
}
......@@ -1316,10 +1316,10 @@ namespace cv { namespace gpu { namespace imgproc
grid.y = divUp(src.rows, threads.y);
if (hrange == 180)
imgproc_krnls::HSV2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::HSV2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
else
imgproc_krnls::HSV2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::HSV2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -1338,7 +1338,7 @@ namespace cv { namespace gpu { namespace imgproc
static const int sector_data[][3] =
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );
HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);
}
......@@ -1355,7 +1355,7 @@ namespace cv { namespace gpu { namespace imgproc
static const int sector_data[][3] =
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );
HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);
}
......@@ -1363,7 +1363,7 @@ namespace cv { namespace gpu { namespace imgproc
/////////////////////////////////////// RGB <-> HLS ////////////////////////////////////////
namespace imgproc_krnls
namespace color_krnls
{
template<typename T, int HR> struct RGB2HLSConvertor;
template<int HR> struct RGB2HLSConvertor<float, HR>
......@@ -1534,7 +1534,7 @@ namespace imgproc_krnls
}
}
namespace cv { namespace gpu { namespace imgproc
namespace cv { namespace gpu { namespace color
{
template <typename T, int SRCCN, int DSTCN>
void RGB2HLS_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)
......@@ -1546,10 +1546,10 @@ namespace cv { namespace gpu { namespace imgproc
grid.y = divUp(src.rows, threads.y);
if (hrange == 180)
imgproc_krnls::RGB2HLS<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2HLS<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
else
imgproc_krnls::RGB2HLS<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::RGB2HLS<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -1591,10 +1591,10 @@ namespace cv { namespace gpu { namespace imgproc
grid.y = divUp(src.rows, threads.y);
if (hrange == 180)
imgproc_krnls::HLS2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::HLS2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
else
imgproc_krnls::HLS2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
color_krnls::HLS2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step,
dst.ptr, dst.step, src.rows, src.cols, bidx);
if (stream == 0)
......@@ -1613,7 +1613,7 @@ namespace cv { namespace gpu { namespace imgproc
static const int sector_data[][3]=
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );
HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);
}
......@@ -1630,7 +1630,7 @@ namespace cv { namespace gpu { namespace imgproc
static const int sector_data[][3]=
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};
cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );
cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );
HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);
}
......
modules/gpu/src/imgproc_gpu.cpp
View file @ 80c1aecf
......@@ -54,8 +54,6 @@ void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }
void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }
double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; }
void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); }
void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); }
......@@ -73,10 +71,8 @@ void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*) { throw_nogpu();
#else /* !defined (HAVE_CUDA) */
namespace cv { namespace gpu
namespace cv { namespace gpu { namespace imgproc
{
namespace imgproc
{
void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
......@@ -88,53 +84,7 @@ namespace cv { namespace gpu
void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);
void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);
void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream);
void Gray2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream);
void RGB2Gray_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB2Gray_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream);
void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);
void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HSV_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HSV2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void RGB2HLS_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
void HLS2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);
}
}}
}}}
////////////////////////////////////////////////////////////////////////
// remap
......@@ -278,375 +228,6 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q,
reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, StreamAccessor::getStream(stream));
}
////////////////////////////////////////////////////////////////////////
// cvtColor
namespace
{
#undef R2Y
#undef G2Y
#undef B2Y
enum
{
yuv_shift = 14,
xyz_shift = 12,
R2Y = 4899,
G2Y = 9617,
B2Y = 1868,
BLOCK_SIZE = 256
};
}
namespace
{
void cvtColor_caller(const GpuMat& src, GpuMat& dst, int code, int dcn, const cudaStream_t& stream)
{
Size sz = src.size();
int scn = src.channels(), depth = src.depth(), bidx;
CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F);
switch (code)
{
case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);
static const func_t funcs[] = {imgproc::RGB2RGB_gpu_8u, 0, imgproc::RGB2RGB_gpu_16u, 0, 0, imgproc::RGB2RGB_gpu_32f};
CV_Assert(scn == 3 || scn == 4);
dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3;
bidx = code == CV_BGR2BGRA || code == CV_BGRA2BGR ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, dcn));
funcs[depth](src, scn, dst, dcn, bidx, stream);
break;
}
case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
{
CV_Assert((scn == 3 || scn == 4) && depth == CV_8U);
int green_bits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565
|| code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;
bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555
|| code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2;
dst.create(sz, CV_8UC2);
imgproc::RGB2RGB5x5_gpu(src, scn, dst, green_bits, bidx, stream);
break;
}
case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
{
if (dcn <= 0) dcn = 3;
CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);
int green_bits = code == CV_BGR5652BGR || code == CV_BGR5652RGB
|| code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;
bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR
|| code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, dcn));
imgproc::RGB5x52RGB_gpu(src, green_bits, dst, dcn, bidx, stream);
break;
}
case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);
static const func_t funcs[] = {imgproc::RGB2Gray_gpu_8u, 0, imgproc::RGB2Gray_gpu_16u, 0, 0, imgproc::RGB2Gray_gpu_32f};
CV_Assert(scn == 3 || scn == 4);
bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
dst.create(sz, CV_MAKETYPE(depth, 1));
funcs[depth](src, scn, dst, bidx, stream);
break;
}
case CV_BGR5652GRAY: case CV_BGR5552GRAY:
{
CV_Assert(scn == 2 && depth == CV_8U);
int green_bits = code == CV_BGR5652GRAY ? 6 : 5;
dst.create(sz, CV_8UC1);
imgproc::RGB5x52Gray_gpu(src, green_bits, dst, stream);
break;
}
case CV_GRAY2BGR: case CV_GRAY2BGRA:
{
typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);
static const func_t funcs[] = {imgproc::Gray2RGB_gpu_8u, 0, imgproc::Gray2RGB_gpu_16u, 0, 0, imgproc::Gray2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert(scn == 1 && (dcn == 3 || dcn == 4));
dst.create(sz, CV_MAKETYPE(depth, dcn));
funcs[depth](src, dst, dcn, stream);
break;
}
case CV_GRAY2BGR565: case CV_GRAY2BGR555:
{
CV_Assert(scn == 1 && depth == CV_8U);
int green_bits = code == CV_GRAY2BGR565 ? 6 : 5;
dst.create(sz, CV_8UC2);
imgproc::Gray2RGB5x5_gpu(src, dst, green_bits, stream);
break;
}
case CV_BGR2YCrCb: case CV_RGB2YCrCb:
case CV_BGR2YUV: case CV_RGB2YUV:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {imgproc::RGB2YCrCb_gpu_8u, 0, imgproc::RGB2YCrCb_gpu_16u, 0, 0, imgproc::RGB2YCrCb_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_BGR2YCrCb || code == CV_RGB2YUV ? 0 : 2;
static const float yuv_f[] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
static const int yuv_i[] = { B2Y, G2Y, R2Y, 8061, 14369 };
static const float YCrCb_f[] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
static const int YCrCb_i[] = {R2Y, G2Y, B2Y, 11682, 9241};
float coeffs_f[5];
int coeffs_i[5];
::memcpy(coeffs_f, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_f : yuv_f, sizeof(yuv_f));
::memcpy(coeffs_i, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_i : yuv_i, sizeof(yuv_i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[2]);
std::swap(coeffs_i[0], coeffs_i[2]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);
break;
}
case CV_YCrCb2BGR: case CV_YCrCb2RGB:
case CV_YUV2BGR: case CV_YUV2RGB:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {imgproc::YCrCb2RGB_gpu_8u, 0, imgproc::YCrCb2RGB_gpu_16u, 0, 0, imgproc::YCrCb2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_YCrCb2BGR || code == CV_YUV2RGB ? 0 : 2;
static const float yuv_f[] = { 2.032f, -0.395f, -0.581f, 1.140f };
static const int yuv_i[] = { 33292, -6472, -9519, 18678 };
static const float YCrCb_f[] = {1.403f, -0.714f, -0.344f, 1.773f};
static const int YCrCb_i[] = {22987, -11698, -5636, 29049};
const float* coeffs_f = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_f : yuv_f;
const int* coeffs_i = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_i : yuv_i;
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);
break;
}
case CV_BGR2XYZ: case CV_RGB2XYZ:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {imgproc::RGB2XYZ_gpu_8u, 0, imgproc::RGB2XYZ_gpu_16u, 0, 0, imgproc::RGB2XYZ_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_BGR2XYZ ? 0 : 2;
static const float RGB2XYZ_D65f[] =
{
0.412453f, 0.357580f, 0.180423f,
0.212671f, 0.715160f, 0.072169f,
0.019334f, 0.119193f, 0.950227f
};
static const int RGB2XYZ_D65i[] =
{
1689, 1465, 739,
871, 2929, 296,
79, 488, 3892
};
float coeffs_f[9];
int coeffs_i[9];
::memcpy(coeffs_f, RGB2XYZ_D65f, sizeof(RGB2XYZ_D65f));
::memcpy(coeffs_i, RGB2XYZ_D65i, sizeof(RGB2XYZ_D65i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[2]);
std::swap(coeffs_f[3], coeffs_f[5]);
std::swap(coeffs_f[6], coeffs_f[8]);
std::swap(coeffs_i[0], coeffs_i[2]);
std::swap(coeffs_i[3], coeffs_i[5]);
std::swap(coeffs_i[6], coeffs_i[8]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, coeffs, stream);
break;
}
case CV_XYZ2BGR: case CV_XYZ2RGB:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn,
const void* coeffs, cudaStream_t stream);
static const func_t funcs[] = {imgproc::XYZ2RGB_gpu_8u, 0, imgproc::XYZ2RGB_gpu_16u, 0, 0, imgproc::XYZ2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));
bidx = code == CV_XYZ2BGR ? 0 : 2;
static const float XYZ2sRGB_D65f[] =
{
3.240479f, -1.53715f, -0.498535f,
-0.969256f, 1.875991f, 0.041556f,
0.055648f, -0.204043f, 1.057311f
};
static const int XYZ2sRGB_D65i[] =
{
13273, -6296, -2042,
-3970, 7684, 170,
228, -836, 4331
};
float coeffs_f[9];
int coeffs_i[9];
::memcpy(coeffs_f, XYZ2sRGB_D65f, sizeof(XYZ2sRGB_D65f));
::memcpy(coeffs_i, XYZ2sRGB_D65i, sizeof(XYZ2sRGB_D65i));
if (bidx == 0)
{
std::swap(coeffs_f[0], coeffs_f[6]);
std::swap(coeffs_f[1], coeffs_f[7]);
std::swap(coeffs_f[2], coeffs_f[8]);
std::swap(coeffs_i[0], coeffs_i[6]);
std::swap(coeffs_i[1], coeffs_i[7]);
std::swap(coeffs_i[2], coeffs_i[8]);
}
dst.create(sz, CV_MAKETYPE(depth, dcn));
const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;
funcs[depth](src, scn, dst, dcn, coeffs_i, stream);
break;
}
case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:
case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
int hrange, cudaStream_t stream);
static const func_t funcs_hsv[] = {imgproc::RGB2HSV_gpu_8u, 0, 0, 0, 0, imgproc::RGB2HSV_gpu_32f};
static const func_t funcs_hls[] = {imgproc::RGB2HLS_gpu_8u, 0, 0, 0, 0, imgproc::RGB2HLS_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == CV_BGR2HSV || code == CV_BGR2HLS ||
code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV ||
code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 255;
dst.create(sz, CV_MAKETYPE(depth, dcn));
if (code == CV_BGR2HSV || code == CV_RGB2HSV || code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL)
funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);
else
funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);
break;
}
case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:
case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:
{
typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx,
int hrange, cudaStream_t stream);
static const func_t funcs_hsv[] = {imgproc::HSV2RGB_gpu_8u, 0, 0, 0, 0, imgproc::HSV2RGB_gpu_32f};
static const func_t funcs_hls[] = {imgproc::HLS2RGB_gpu_8u, 0, 0, 0, 0, imgproc::HLS2RGB_gpu_32f};
if (dcn <= 0) dcn = 3;
CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));
bidx = code == CV_HSV2BGR || code == CV_HLS2BGR ||
code == CV_HSV2BGR_FULL || code == CV_HLS2BGR_FULL ? 0 : 2;
int hrange = depth == CV_32F ? 360 : code == CV_HSV2BGR || code == CV_HSV2RGB ||
code == CV_HLS2BGR || code == CV_HLS2RGB ? 180 : 255;
dst.create(sz, CV_MAKETYPE(depth, dcn));
if (code == CV_HSV2BGR || code == CV_HSV2RGB || code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL)
funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);
else
funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);
break;
}
default:
CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );
}
}
}
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)
{
cvtColor_caller(src, dst, code, dcn, 0);
}
void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)
{
cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));
}
////////////////////////////////////////////////////////////////////////
// threshold
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment