From 80c1aecfe58fc000fd21395c9f1edcf411fc6835 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Tue, 26 Oct 2010 09:37:24 +0000 Subject: [PATCH] moved gpu::cvtColor from imgproc_gpu.cpp to color.cpp --- modules/gpu/src/color.cpp | 469 ++++++++++++++++++++++++++++++++ modules/gpu/src/cuda/color.cu | 106 ++++---- modules/gpu/src/imgproc_gpu.cpp | 441 +----------------------------- 3 files changed, 533 insertions(+), 483 deletions(-) create mode 100644 modules/gpu/src/color.cpp diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp new file mode 100644 index 0000000000..e6be16f265 --- /dev/null +++ b/modules/gpu/src/color.cpp @@ -0,0 +1,469 @@ +/*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) */ diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index e956ff8dd7..f368dd2132 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -54,7 +54,7 @@ using namespace cv::gpu; #define FLT_EPSILON 1.192092896e-07F #endif -namespace imgproc_krnls +namespace color_krnls { template struct ColorChannel {}; template<> struct ColorChannel @@ -99,7 +99,7 @@ namespace imgproc_krnls ////////////////// Various 3/4-channel to 3/4-channel RGB transformations ///////////////// -namespace imgproc_krnls +namespace color_krnls { template __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 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<<>>(src.ptr, src.step, + color_krnls::RGB2RGB<<>>(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 struct RGB5x52RGBConverter {}; template struct RGB5x52RGBConverter<5, DSTCN> @@ -274,7 +274,7 @@ namespace imgproc_krnls } } -namespace cv { namespace gpu { namespace imgproc +namespace cv { namespace gpu { namespace color { template 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<<>>(src.ptr, src.step, + color_krnls::RGB5x52RGB<<>>(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<<>>(src.ptr, src.step, + color_krnls::RGB2RGB5x5<<>>(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 __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 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<<>>(src.ptr, src.step, + color_krnls::Gray2RGB<<>>(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<<>>(src.ptr, src.step, + color_krnls::Gray2RGB5x5<<>>(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 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<<>>(src.ptr, src.step, + color_krnls::RGB2Gray<<>>(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<<>>(src.ptr, src.step, + color_krnls::RGB5x52Gray<<>>(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 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<<>>(src.ptr, src.step, + color_krnls::RGB2YCrCb<<>>(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, RGB2YCrCb_caller} }; - 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, RGB2YCrCb_caller} }; - 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, RGB2YCrCb_caller} }; - 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<<>>(src.ptr, src.step, + color_krnls::YCrCb2RGB<<>>(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, YCrCb2RGB_caller} }; - 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, YCrCb2RGB_caller} }; - 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, YCrCb2RGB_caller} }; - 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 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<<>>(src.ptr, src.step, + color_krnls::RGB2XYZ<<>>(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, RGB2XYZ_caller} }; - 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, RGB2XYZ_caller} }; - 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, RGB2XYZ_caller} }; - 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<<>>(src.ptr, src.step, + color_krnls::XYZ2RGB<<>>(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, XYZ2RGB_caller} }; - 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, XYZ2RGB_caller} }; - 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, XYZ2RGB_caller} }; - 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 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<<>>(src.ptr, src.step, + color_krnls::RGB2HSV<<>>(src.ptr, src.step, dst.ptr, dst.step, src.rows, src.cols, bidx); else - imgproc_krnls::RGB2HSV<<>>(src.ptr, src.step, + color_krnls::RGB2HSV<<>>(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<<>>(src.ptr, src.step, + color_krnls::HSV2RGB<<>>(src.ptr, src.step, dst.ptr, dst.step, src.rows, src.cols, bidx); else - imgproc_krnls::HSV2RGB<<>>(src.ptr, src.step, + color_krnls::HSV2RGB<<>>(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 struct RGB2HLSConvertor; template struct RGB2HLSConvertor @@ -1534,7 +1534,7 @@ namespace imgproc_krnls } } -namespace cv { namespace gpu { namespace imgproc +namespace cv { namespace gpu { namespace color { template 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<<>>(src.ptr, src.step, + color_krnls::RGB2HLS<<>>(src.ptr, src.step, dst.ptr, dst.step, src.rows, src.cols, bidx); else - imgproc_krnls::RGB2HLS<<>>(src.ptr, src.step, + color_krnls::RGB2HLS<<>>(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<<>>(src.ptr, src.step, + color_krnls::HLS2RGB<<>>(src.ptr, src.step, dst.ptr, dst.step, src.rows, src.cols, bidx); else - imgproc_krnls::HLS2RGB<<>>(src.ptr, src.step, + color_krnls::HLS2RGB<<>>(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); } diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index b14198aa2a..c7619120a5 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -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,68 +71,20 @@ void cv::gpu::histRange(const GpuMat&, GpuMat*, const GpuMat*) { throw_nogpu(); #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace gpu -{ - 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); - - extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps); - extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps); - - void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream); - void drawColorDisp_gpu(const DevMem2D_& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream); - - void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); - void reprojectImageTo3D_gpu(const DevMem2D_& 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); +namespace cv { namespace gpu { 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); - 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); + extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps); + extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps); - 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 drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream); + void drawColorDisp_gpu(const DevMem2D_& src, const DevMem2D& dst, int ndisp, const 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); - } -}} + void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); + void reprojectImageTo3D_gpu(const DevMem2D_& disp, const DevMem2Df& xyzw, const float* q, const 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 -- GitLab