--- /dev/null
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of the copyright holders may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include "precomp.hpp"\r
+\r
+using namespace cv;\r
+using namespace cv::gpu;\r
+\r
+#if !defined (HAVE_CUDA)\r
+\r
+void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }\r
+void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }\r
+\r
+#else /* !defined (HAVE_CUDA) */\r
+\r
+namespace cv { namespace gpu { namespace color \r
+{\r
+ void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+\r
+ void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream);\r
+\r
+ void Gray2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
+ void Gray2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
+ void Gray2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
+ void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream);\r
+\r
+ void RGB2Gray_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ void RGB2Gray_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+ void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+ void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+ void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+\r
+ void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+ void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+ void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
+\r
+ void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+ void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+ void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+\r
+ void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+ void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+ void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
+\r
+ void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ void RGB2HSV_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+\r
+ void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ void HSV2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+\r
+ void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ void RGB2HLS_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+\r
+ void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ void HLS2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+}}}\r
+\r
+namespace\r
+{\r
+ #undef R2Y\r
+ #undef G2Y\r
+ #undef B2Y\r
+ \r
+ enum\r
+ {\r
+ yuv_shift = 14,\r
+ xyz_shift = 12,\r
+ R2Y = 4899,\r
+ G2Y = 9617,\r
+ B2Y = 1868,\r
+ BLOCK_SIZE = 256\r
+ };\r
+}\r
+\r
+namespace\r
+{\r
+ void cvtColor_caller(const GpuMat& src, GpuMat& dst, int code, int dcn, const cudaStream_t& stream) \r
+ {\r
+ Size sz = src.size();\r
+ int scn = src.channels(), depth = src.depth(), bidx;\r
+ \r
+ CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F);\r
+\r
+ switch (code)\r
+ {\r
+ case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:\r
+ case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: \r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::RGB2RGB_gpu_8u, 0, color::RGB2RGB_gpu_16u, 0, 0, color::RGB2RGB_gpu_32f};\r
+\r
+ CV_Assert(scn == 3 || scn == 4);\r
+\r
+ dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3;\r
+ bidx = code == CV_BGR2BGRA || code == CV_BGRA2BGR ? 0 : 2;\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ funcs[depth](src, scn, dst, dcn, bidx, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:\r
+ case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:\r
+ {\r
+ CV_Assert((scn == 3 || scn == 4) && depth == CV_8U);\r
+\r
+ int green_bits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 \r
+ || code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;\r
+ bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555 \r
+ || code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2;\r
+\r
+ dst.create(sz, CV_8UC2);\r
+\r
+ color::RGB2RGB5x5_gpu(src, scn, dst, green_bits, bidx, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:\r
+ case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:\r
+ {\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);\r
+\r
+ int green_bits = code == CV_BGR5652BGR || code == CV_BGR5652RGB \r
+ || code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;\r
+ bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR \r
+ || code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2;\r
+\r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ color::RGB5x52RGB_gpu(src, green_bits, dst, dcn, bidx, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::RGB2Gray_gpu_8u, 0, color::RGB2Gray_gpu_16u, 0, 0, color::RGB2Gray_gpu_32f};\r
+\r
+ CV_Assert(scn == 3 || scn == 4);\r
+ \r
+ bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;\r
+\r
+ dst.create(sz, CV_MAKETYPE(depth, 1));\r
+\r
+ funcs[depth](src, scn, dst, bidx, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_BGR5652GRAY: case CV_BGR5552GRAY:\r
+ {\r
+ CV_Assert(scn == 2 && depth == CV_8U);\r
+\r
+ int green_bits = code == CV_BGR5652GRAY ? 6 : 5;\r
+\r
+ dst.create(sz, CV_8UC1);\r
+\r
+ color::RGB5x52Gray_gpu(src, green_bits, dst, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_GRAY2BGR: case CV_GRAY2BGRA:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::Gray2RGB_gpu_8u, 0, color::Gray2RGB_gpu_16u, 0, 0, color::Gray2RGB_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert(scn == 1 && (dcn == 3 || dcn == 4));\r
+\r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ funcs[depth](src, dst, dcn, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_GRAY2BGR565: case CV_GRAY2BGR555:\r
+ {\r
+ CV_Assert(scn == 1 && depth == CV_8U);\r
+\r
+ int green_bits = code == CV_GRAY2BGR565 ? 6 : 5;\r
+\r
+ dst.create(sz, CV_8UC2);\r
+ \r
+ color::Gray2RGB5x5_gpu(src, dst, green_bits, stream);\r
+ break;\r
+ }\r
+\r
+ case CV_BGR2YCrCb: case CV_RGB2YCrCb:\r
+ case CV_BGR2YUV: case CV_RGB2YUV:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
+ const void* coeffs, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::RGB2YCrCb_gpu_8u, 0, color::RGB2YCrCb_gpu_16u, 0, 0, color::RGB2YCrCb_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
+\r
+ bidx = code == CV_BGR2YCrCb || code == CV_RGB2YUV ? 0 : 2;\r
+\r
+ static const float yuv_f[] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };\r
+ static const int yuv_i[] = { B2Y, G2Y, R2Y, 8061, 14369 };\r
+\r
+ static const float YCrCb_f[] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};\r
+ static const int YCrCb_i[] = {R2Y, G2Y, B2Y, 11682, 9241};\r
+\r
+ float coeffs_f[5];\r
+ int coeffs_i[5];\r
+ ::memcpy(coeffs_f, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_f : yuv_f, sizeof(yuv_f));\r
+ ::memcpy(coeffs_i, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_i : yuv_i, sizeof(yuv_i));\r
+\r
+ if (bidx == 0) \r
+ {\r
+ std::swap(coeffs_f[0], coeffs_f[2]);\r
+ std::swap(coeffs_i[0], coeffs_i[2]);\r
+ }\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
+\r
+ funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_YCrCb2BGR: case CV_YCrCb2RGB:\r
+ case CV_YUV2BGR: case CV_YUV2RGB:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
+ const void* coeffs, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::YCrCb2RGB_gpu_8u, 0, color::YCrCb2RGB_gpu_16u, 0, 0, color::YCrCb2RGB_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
+\r
+ bidx = code == CV_YCrCb2BGR || code == CV_YUV2RGB ? 0 : 2;\r
+\r
+ static const float yuv_f[] = { 2.032f, -0.395f, -0.581f, 1.140f };\r
+ static const int yuv_i[] = { 33292, -6472, -9519, 18678 }; \r
+\r
+ static const float YCrCb_f[] = {1.403f, -0.714f, -0.344f, 1.773f};\r
+ static const int YCrCb_i[] = {22987, -11698, -5636, 29049};\r
+\r
+ const float* coeffs_f = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_f : yuv_f;\r
+ const int* coeffs_i = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_i : yuv_i;\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+ \r
+ const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
+\r
+ funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_BGR2XYZ: case CV_RGB2XYZ:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, \r
+ const void* coeffs, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::RGB2XYZ_gpu_8u, 0, color::RGB2XYZ_gpu_16u, 0, 0, color::RGB2XYZ_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
+\r
+ bidx = code == CV_BGR2XYZ ? 0 : 2;\r
+\r
+ static const float RGB2XYZ_D65f[] =\r
+ {\r
+ 0.412453f, 0.357580f, 0.180423f,\r
+ 0.212671f, 0.715160f, 0.072169f,\r
+ 0.019334f, 0.119193f, 0.950227f\r
+ };\r
+ static const int RGB2XYZ_D65i[] =\r
+ {\r
+ 1689, 1465, 739,\r
+ 871, 2929, 296,\r
+ 79, 488, 3892\r
+ };\r
+\r
+ float coeffs_f[9];\r
+ int coeffs_i[9];\r
+ ::memcpy(coeffs_f, RGB2XYZ_D65f, sizeof(RGB2XYZ_D65f));\r
+ ::memcpy(coeffs_i, RGB2XYZ_D65i, sizeof(RGB2XYZ_D65i));\r
+\r
+ if (bidx == 0) \r
+ {\r
+ std::swap(coeffs_f[0], coeffs_f[2]);\r
+ std::swap(coeffs_f[3], coeffs_f[5]);\r
+ std::swap(coeffs_f[6], coeffs_f[8]);\r
+ \r
+ std::swap(coeffs_i[0], coeffs_i[2]);\r
+ std::swap(coeffs_i[3], coeffs_i[5]);\r
+ std::swap(coeffs_i[6], coeffs_i[8]);\r
+ }\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+ \r
+ const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
+ \r
+ funcs[depth](src, scn, dst, dcn, coeffs, stream);\r
+ break;\r
+ }\r
+ \r
+ case CV_XYZ2BGR: case CV_XYZ2RGB:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, \r
+ const void* coeffs, cudaStream_t stream);\r
+ static const func_t funcs[] = {color::XYZ2RGB_gpu_8u, 0, color::XYZ2RGB_gpu_16u, 0, 0, color::XYZ2RGB_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
+\r
+ bidx = code == CV_XYZ2BGR ? 0 : 2;\r
+\r
+ static const float XYZ2sRGB_D65f[] =\r
+ {\r
+ 3.240479f, -1.53715f, -0.498535f,\r
+ -0.969256f, 1.875991f, 0.041556f,\r
+ 0.055648f, -0.204043f, 1.057311f\r
+ };\r
+ static const int XYZ2sRGB_D65i[] =\r
+ {\r
+ 13273, -6296, -2042,\r
+ -3970, 7684, 170,\r
+ 228, -836, 4331\r
+ };\r
+\r
+ float coeffs_f[9];\r
+ int coeffs_i[9];\r
+ ::memcpy(coeffs_f, XYZ2sRGB_D65f, sizeof(XYZ2sRGB_D65f));\r
+ ::memcpy(coeffs_i, XYZ2sRGB_D65i, sizeof(XYZ2sRGB_D65i));\r
+\r
+ if (bidx == 0) \r
+ {\r
+ std::swap(coeffs_f[0], coeffs_f[6]);\r
+ std::swap(coeffs_f[1], coeffs_f[7]);\r
+ std::swap(coeffs_f[2], coeffs_f[8]);\r
+ \r
+ std::swap(coeffs_i[0], coeffs_i[6]);\r
+ std::swap(coeffs_i[1], coeffs_i[7]);\r
+ std::swap(coeffs_i[2], coeffs_i[8]);\r
+ }\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+ \r
+ const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
+\r
+ funcs[depth](src, scn, dst, dcn, coeffs_i, stream);\r
+ break;\r
+ }\r
+\r
+ case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:\r
+ case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
+ int hrange, cudaStream_t stream);\r
+ static const func_t funcs_hsv[] = {color::RGB2HSV_gpu_8u, 0, 0, 0, 0, color::RGB2HSV_gpu_32f};\r
+ static const func_t funcs_hls[] = {color::RGB2HLS_gpu_8u, 0, 0, 0, 0, color::RGB2HLS_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));\r
+\r
+ bidx = code == CV_BGR2HSV || code == CV_BGR2HLS ||\r
+ code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2;\r
+ int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV ||\r
+ code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 255;\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ if (code == CV_BGR2HSV || code == CV_RGB2HSV || code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL) \r
+ funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
+ else\r
+ funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
+ break;\r
+ }\r
+\r
+ case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:\r
+ case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:\r
+ {\r
+ typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
+ int hrange, cudaStream_t stream);\r
+ static const func_t funcs_hsv[] = {color::HSV2RGB_gpu_8u, 0, 0, 0, 0, color::HSV2RGB_gpu_32f};\r
+ static const func_t funcs_hls[] = {color::HLS2RGB_gpu_8u, 0, 0, 0, 0, color::HLS2RGB_gpu_32f};\r
+\r
+ if (dcn <= 0) dcn = 3;\r
+\r
+ CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));\r
+\r
+ bidx = code == CV_HSV2BGR || code == CV_HLS2BGR ||\r
+ code == CV_HSV2BGR_FULL || code == CV_HLS2BGR_FULL ? 0 : 2;\r
+ int hrange = depth == CV_32F ? 360 : code == CV_HSV2BGR || code == CV_HSV2RGB ||\r
+ code == CV_HLS2BGR || code == CV_HLS2RGB ? 180 : 255;\r
+ \r
+ dst.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ if (code == CV_HSV2BGR || code == CV_HSV2RGB || code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL)\r
+ funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
+ else\r
+ funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
+ break;\r
+ }\r
+\r
+ default:\r
+ CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
+ }\r
+ }\r
+}\r
+\r
+void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)\r
+{\r
+ cvtColor_caller(src, dst, code, dcn, 0);\r
+}\r
+\r
+void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)\r
+{\r
+ cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));\r
+}\r
+\r
+#endif /* !defined (HAVE_CUDA) */\r
#define FLT_EPSILON 1.192092896e-07F\r
#endif\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
template<typename T> struct ColorChannel {};\r
template<> struct ColorChannel<uchar>\r
\r
////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
template <int SRCCN, int DSTCN, typename T>\r
__global__ void RGB2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN, int DSTCN>\r
void RGB2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
\r
/////////// Transforming 16-bit (565 or 555) RGB to/from 24/32-bit (888[8]) RGB //////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
template <int GREEN_BITS, int DSTCN> struct RGB5x52RGBConverter {}; \r
template <int DSTCN> struct RGB5x52RGBConverter<5, DSTCN>\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <int GREEN_BITS, int DSTCN>\r
void RGB5x52RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB5x52RGB<GREEN_BITS, DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB5x52RGB<GREEN_BITS, DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB2RGB5x5<SRCCN, GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2RGB5x5<SRCCN, GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
\r
///////////////////////////////// Grayscale to Color ////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
template <int DSTCN, typename T>\r
__global__ void Gray2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols)\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int DSTCN>\r
void Gray2RGB_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::Gray2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::Gray2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::Gray2RGB5x5<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::Gray2RGB5x5<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
\r
///////////////////////////////// Color to Grayscale ////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
#undef R2Y\r
#undef G2Y\r
} \r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN>\r
void RGB2Gray_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB2Gray<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2Gray<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB5x52Gray<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB5x52Gray<GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
\r
///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
__constant__ float cYCrCbCoeffs_f[5];\r
__constant__ int cYCrCbCoeffs_i[5];\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN, int DSTCN>\r
void RGB2YCrCb_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB2YCrCb<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2YCrCb<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
{RGB2YCrCb_caller<uchar, 4, 3>, RGB2YCrCb_caller<uchar, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
\r
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
{RGB2YCrCb_caller<ushort, 4, 3>, RGB2YCrCb_caller<ushort, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
\r
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
{RGB2YCrCb_caller<float, 4, 3>, RGB2YCrCb_caller<float, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) );\r
\r
RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::YCrCb2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::YCrCb2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
{YCrCb2RGB_caller<uchar, 4, 3>, YCrCb2RGB_caller<uchar, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
\r
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
{YCrCb2RGB_caller<ushort, 4, 3>, YCrCb2RGB_caller<ushort, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
\r
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
{YCrCb2RGB_caller<float, 4, 3>, YCrCb2RGB_caller<float, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) );\r
\r
YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
\r
////////////////////////////////////// RGB <-> XYZ ///////////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
__constant__ float cXYZ_D65f[9];\r
__constant__ int cXYZ_D65i[9];\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN, int DSTCN>\r
void RGB2XYZ_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::RGB2XYZ<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2XYZ<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
{RGB2XYZ_caller<uchar, 4, 3>, RGB2XYZ_caller<uchar, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
\r
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
{RGB2XYZ_caller<ushort, 4, 3>, RGB2XYZ_caller<ushort, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
\r
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
{RGB2XYZ_caller<float, 4, 3>, RGB2XYZ_caller<float, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
\r
RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- imgproc_krnls::XYZ2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::XYZ2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
{XYZ2RGB_caller<uchar, 4, 3>, XYZ2RGB_caller<uchar, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
\r
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
{XYZ2RGB_caller<ushort, 4, 3>, XYZ2RGB_caller<ushort, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
\r
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
{XYZ2RGB_caller<float, 4, 3>, XYZ2RGB_caller<float, 4, 4>}\r
};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
\r
XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
}\r
\r
////////////////////////////////////// RGB <-> HSV ///////////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
__constant__ int cHsvDivTable[256];\r
\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN, int DSTCN>\r
void RGB2HSV_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)\r
grid.y = divUp(src.rows, threads.y);\r
\r
if (hrange == 180)\r
- imgproc_krnls::RGB2HSV<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2HSV<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
else\r
- imgproc_krnls::RGB2HSV<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2HSV<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
4352, 4334, 4316, 4298, 4281, 4263, 4246, 4229,\r
4212, 4195, 4178, 4161, 4145, 4128, 4112, 4096\r
};\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvDivTable, div_table, sizeof(div_table)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvDivTable, div_table, sizeof(div_table)) );\r
\r
RGB2HSV_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);\r
}\r
grid.y = divUp(src.rows, threads.y);\r
\r
if (hrange == 180)\r
- imgproc_krnls::HSV2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::HSV2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
else\r
- imgproc_krnls::HSV2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::HSV2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
static const int sector_data[][3] =\r
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );\r
\r
HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);\r
}\r
static const int sector_data[][3] =\r
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHsvSectorData, sector_data, sizeof(sector_data)) );\r
\r
HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);\r
}\r
\r
/////////////////////////////////////// RGB <-> HLS ////////////////////////////////////////\r
\r
-namespace imgproc_krnls\r
+namespace color_krnls\r
{\r
template<typename T, int HR> struct RGB2HLSConvertor;\r
template<int HR> struct RGB2HLSConvertor<float, HR>\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace imgproc\r
+namespace cv { namespace gpu { namespace color\r
{\r
template <typename T, int SRCCN, int DSTCN>\r
void RGB2HLS_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream)\r
grid.y = divUp(src.rows, threads.y);\r
\r
if (hrange == 180)\r
- imgproc_krnls::RGB2HLS<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2HLS<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
else\r
- imgproc_krnls::RGB2HLS<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::RGB2HLS<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
grid.y = divUp(src.rows, threads.y);\r
\r
if (hrange == 180)\r
- imgproc_krnls::HLS2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::HLS2RGB<SRCCN, DSTCN, 180, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
else\r
- imgproc_krnls::HLS2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ color_krnls::HLS2RGB<SRCCN, DSTCN, 255, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
static const int sector_data[][3]=\r
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );\r
\r
HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);\r
}\r
static const int sector_data[][3]=\r
{{1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0}};\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(color_krnls::cHlsSectorData, sector_data, sizeof(sector_data)) );\r
\r
HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream);\r
}\r
void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }\r
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }\r
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, const Stream&) { throw_nogpu(); }\r
-void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int) { throw_nogpu(); }\r
-void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, const Stream&) { throw_nogpu(); }\r
double cv::gpu::threshold(const GpuMat&, GpuMat&, double) { throw_nogpu(); return 0.0; }\r
void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int) { throw_nogpu(); }\r
void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&) { throw_nogpu(); }\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
-namespace cv { namespace gpu \r
-{ \r
- namespace imgproc \r
- {\r
- void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);\r
- void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);\r
-\r
- extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps);\r
- extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps);\r
-\r
- void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
- void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
-\r
- void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
- void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
-\r
- void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
- void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
- void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
-\r
- void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
- void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream);\r
-\r
- void Gray2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
- void Gray2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
- void Gray2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
- void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream);\r
-\r
- void RGB2Gray_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
- void RGB2Gray_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
- void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
- void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream);\r
-\r
- void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
- void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
- void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
-\r
- void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
- void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
- void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream);\r
-\r
- void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
- void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
- void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
-\r
- void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
- void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
- void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream);\r
-\r
- void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
- void RGB2HSV_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+namespace cv { namespace gpu { namespace imgproc \r
+{\r
+ void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);\r
+ void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);\r
\r
- void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
- void HSV2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps);\r
+ extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps);\r
\r
- void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
- void RGB2HLS_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
+ void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
+ void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
\r
- void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
- void HLS2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream);\r
- }\r
-}}\r
+ void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
+ void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); \r
+}}}\r
\r
////////////////////////////////////////////////////////////////////////\r
// remap\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
-// cvtColor\r
-\r
-namespace\r
-{\r
- #undef R2Y\r
- #undef G2Y\r
- #undef B2Y\r
- \r
- enum\r
- {\r
- yuv_shift = 14,\r
- xyz_shift = 12,\r
- R2Y = 4899,\r
- G2Y = 9617,\r
- B2Y = 1868,\r
- BLOCK_SIZE = 256\r
- };\r
-}\r
-\r
-namespace\r
-{\r
- void cvtColor_caller(const GpuMat& src, GpuMat& dst, int code, int dcn, const cudaStream_t& stream) \r
- {\r
- Size sz = src.size();\r
- int scn = src.channels(), depth = src.depth(), bidx;\r
- \r
- CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F);\r
-\r
- switch (code)\r
- {\r
- case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:\r
- case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: \r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::RGB2RGB_gpu_8u, 0, imgproc::RGB2RGB_gpu_16u, 0, 0, imgproc::RGB2RGB_gpu_32f};\r
-\r
- CV_Assert(scn == 3 || scn == 4);\r
-\r
- dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3;\r
- bidx = code == CV_BGR2BGRA || code == CV_BGRA2BGR ? 0 : 2;\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- funcs[depth](src, scn, dst, dcn, bidx, stream);\r
- break;\r
- }\r
- \r
- case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:\r
- case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:\r
- {\r
- CV_Assert((scn == 3 || scn == 4) && depth == CV_8U);\r
-\r
- int green_bits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 \r
- || code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;\r
- bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555 \r
- || code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2;\r
-\r
- dst.create(sz, CV_8UC2);\r
-\r
- imgproc::RGB2RGB5x5_gpu(src, scn, dst, green_bits, bidx, stream);\r
- break;\r
- }\r
- \r
- case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:\r
- case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:\r
- {\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);\r
-\r
- int green_bits = code == CV_BGR5652BGR || code == CV_BGR5652RGB \r
- || code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;\r
- bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR \r
- || code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2;\r
-\r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- imgproc::RGB5x52RGB_gpu(src, green_bits, dst, dcn, bidx, stream);\r
- break;\r
- }\r
- \r
- case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::RGB2Gray_gpu_8u, 0, imgproc::RGB2Gray_gpu_16u, 0, 0, imgproc::RGB2Gray_gpu_32f};\r
-\r
- CV_Assert(scn == 3 || scn == 4);\r
- \r
- bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;\r
-\r
- dst.create(sz, CV_MAKETYPE(depth, 1));\r
-\r
- funcs[depth](src, scn, dst, bidx, stream);\r
- break;\r
- }\r
- \r
- case CV_BGR5652GRAY: case CV_BGR5552GRAY:\r
- {\r
- CV_Assert(scn == 2 && depth == CV_8U);\r
-\r
- int green_bits = code == CV_BGR5652GRAY ? 6 : 5;\r
-\r
- dst.create(sz, CV_8UC1);\r
-\r
- imgproc::RGB5x52Gray_gpu(src, green_bits, dst, stream);\r
- break;\r
- }\r
- \r
- case CV_GRAY2BGR: case CV_GRAY2BGRA:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::Gray2RGB_gpu_8u, 0, imgproc::Gray2RGB_gpu_16u, 0, 0, imgproc::Gray2RGB_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert(scn == 1 && (dcn == 3 || dcn == 4));\r
-\r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- funcs[depth](src, dst, dcn, stream);\r
- break;\r
- }\r
- \r
- case CV_GRAY2BGR565: case CV_GRAY2BGR555:\r
- {\r
- CV_Assert(scn == 1 && depth == CV_8U);\r
-\r
- int green_bits = code == CV_GRAY2BGR565 ? 6 : 5;\r
-\r
- dst.create(sz, CV_8UC2);\r
- \r
- imgproc::Gray2RGB5x5_gpu(src, dst, green_bits, stream);\r
- break;\r
- }\r
-\r
- case CV_BGR2YCrCb: case CV_RGB2YCrCb:\r
- case CV_BGR2YUV: case CV_RGB2YUV:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
- const void* coeffs, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::RGB2YCrCb_gpu_8u, 0, imgproc::RGB2YCrCb_gpu_16u, 0, 0, imgproc::RGB2YCrCb_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
-\r
- bidx = code == CV_BGR2YCrCb || code == CV_RGB2YUV ? 0 : 2;\r
-\r
- static const float yuv_f[] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };\r
- static const int yuv_i[] = { B2Y, G2Y, R2Y, 8061, 14369 };\r
-\r
- static const float YCrCb_f[] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};\r
- static const int YCrCb_i[] = {R2Y, G2Y, B2Y, 11682, 9241};\r
-\r
- float coeffs_f[5];\r
- int coeffs_i[5];\r
- ::memcpy(coeffs_f, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_f : yuv_f, sizeof(yuv_f));\r
- ::memcpy(coeffs_i, code == CV_BGR2YCrCb || code == CV_RGB2YCrCb ? YCrCb_i : yuv_i, sizeof(yuv_i));\r
-\r
- if (bidx == 0) \r
- {\r
- std::swap(coeffs_f[0], coeffs_f[2]);\r
- std::swap(coeffs_i[0], coeffs_i[2]);\r
- }\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
-\r
- funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);\r
- break;\r
- }\r
- \r
- case CV_YCrCb2BGR: case CV_YCrCb2RGB:\r
- case CV_YUV2BGR: case CV_YUV2RGB:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
- const void* coeffs, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::YCrCb2RGB_gpu_8u, 0, imgproc::YCrCb2RGB_gpu_16u, 0, 0, imgproc::YCrCb2RGB_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
-\r
- bidx = code == CV_YCrCb2BGR || code == CV_YUV2RGB ? 0 : 2;\r
-\r
- static const float yuv_f[] = { 2.032f, -0.395f, -0.581f, 1.140f };\r
- static const int yuv_i[] = { 33292, -6472, -9519, 18678 }; \r
-\r
- static const float YCrCb_f[] = {1.403f, -0.714f, -0.344f, 1.773f};\r
- static const int YCrCb_i[] = {22987, -11698, -5636, 29049};\r
-\r
- const float* coeffs_f = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_f : yuv_f;\r
- const int* coeffs_i = code == CV_YCrCb2BGR || code == CV_YCrCb2RGB ? YCrCb_i : yuv_i;\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
- \r
- const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
-\r
- funcs[depth](src, scn, dst, dcn, bidx, coeffs, stream);\r
- break;\r
- }\r
- \r
- case CV_BGR2XYZ: case CV_RGB2XYZ:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, \r
- const void* coeffs, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::RGB2XYZ_gpu_8u, 0, imgproc::RGB2XYZ_gpu_16u, 0, 0, imgproc::RGB2XYZ_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
-\r
- bidx = code == CV_BGR2XYZ ? 0 : 2;\r
-\r
- static const float RGB2XYZ_D65f[] =\r
- {\r
- 0.412453f, 0.357580f, 0.180423f,\r
- 0.212671f, 0.715160f, 0.072169f,\r
- 0.019334f, 0.119193f, 0.950227f\r
- };\r
- static const int RGB2XYZ_D65i[] =\r
- {\r
- 1689, 1465, 739,\r
- 871, 2929, 296,\r
- 79, 488, 3892\r
- };\r
-\r
- float coeffs_f[9];\r
- int coeffs_i[9];\r
- ::memcpy(coeffs_f, RGB2XYZ_D65f, sizeof(RGB2XYZ_D65f));\r
- ::memcpy(coeffs_i, RGB2XYZ_D65i, sizeof(RGB2XYZ_D65i));\r
-\r
- if (bidx == 0) \r
- {\r
- std::swap(coeffs_f[0], coeffs_f[2]);\r
- std::swap(coeffs_f[3], coeffs_f[5]);\r
- std::swap(coeffs_f[6], coeffs_f[8]);\r
- \r
- std::swap(coeffs_i[0], coeffs_i[2]);\r
- std::swap(coeffs_i[3], coeffs_i[5]);\r
- std::swap(coeffs_i[6], coeffs_i[8]);\r
- }\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
- \r
- const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
- \r
- funcs[depth](src, scn, dst, dcn, coeffs, stream);\r
- break;\r
- }\r
- \r
- case CV_XYZ2BGR: case CV_XYZ2RGB:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, \r
- const void* coeffs, cudaStream_t stream);\r
- static const func_t funcs[] = {imgproc::XYZ2RGB_gpu_8u, 0, imgproc::XYZ2RGB_gpu_16u, 0, 0, imgproc::XYZ2RGB_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4));\r
-\r
- bidx = code == CV_XYZ2BGR ? 0 : 2;\r
-\r
- static const float XYZ2sRGB_D65f[] =\r
- {\r
- 3.240479f, -1.53715f, -0.498535f,\r
- -0.969256f, 1.875991f, 0.041556f,\r
- 0.055648f, -0.204043f, 1.057311f\r
- };\r
- static const int XYZ2sRGB_D65i[] =\r
- {\r
- 13273, -6296, -2042,\r
- -3970, 7684, 170,\r
- 228, -836, 4331\r
- };\r
-\r
- float coeffs_f[9];\r
- int coeffs_i[9];\r
- ::memcpy(coeffs_f, XYZ2sRGB_D65f, sizeof(XYZ2sRGB_D65f));\r
- ::memcpy(coeffs_i, XYZ2sRGB_D65i, sizeof(XYZ2sRGB_D65i));\r
-\r
- if (bidx == 0) \r
- {\r
- std::swap(coeffs_f[0], coeffs_f[6]);\r
- std::swap(coeffs_f[1], coeffs_f[7]);\r
- std::swap(coeffs_f[2], coeffs_f[8]);\r
- \r
- std::swap(coeffs_i[0], coeffs_i[6]);\r
- std::swap(coeffs_i[1], coeffs_i[7]);\r
- std::swap(coeffs_i[2], coeffs_i[8]);\r
- }\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
- \r
- const void* coeffs = depth == CV_32F ? (void*)coeffs_f : (void*)coeffs_i;\r
-\r
- funcs[depth](src, scn, dst, dcn, coeffs_i, stream);\r
- break;\r
- }\r
-\r
- case CV_BGR2HSV: case CV_RGB2HSV: case CV_BGR2HSV_FULL: case CV_RGB2HSV_FULL:\r
- case CV_BGR2HLS: case CV_RGB2HLS: case CV_BGR2HLS_FULL: case CV_RGB2HLS_FULL:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
- int hrange, cudaStream_t stream);\r
- static const func_t funcs_hsv[] = {imgproc::RGB2HSV_gpu_8u, 0, 0, 0, 0, imgproc::RGB2HSV_gpu_32f};\r
- static const func_t funcs_hls[] = {imgproc::RGB2HLS_gpu_8u, 0, 0, 0, 0, imgproc::RGB2HLS_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));\r
-\r
- bidx = code == CV_BGR2HSV || code == CV_BGR2HLS ||\r
- code == CV_BGR2HSV_FULL || code == CV_BGR2HLS_FULL ? 0 : 2;\r
- int hrange = depth == CV_32F ? 360 : code == CV_BGR2HSV || code == CV_RGB2HSV ||\r
- code == CV_BGR2HLS || code == CV_RGB2HLS ? 180 : 255;\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- if (code == CV_BGR2HSV || code == CV_RGB2HSV || code == CV_BGR2HSV_FULL || code == CV_RGB2HSV_FULL) \r
- funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
- else\r
- funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
- break;\r
- }\r
-\r
- case CV_HSV2BGR: case CV_HSV2RGB: case CV_HSV2BGR_FULL: case CV_HSV2RGB_FULL:\r
- case CV_HLS2BGR: case CV_HLS2RGB: case CV_HLS2BGR_FULL: case CV_HLS2RGB_FULL:\r
- {\r
- typedef void (*func_t)(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, \r
- int hrange, cudaStream_t stream);\r
- static const func_t funcs_hsv[] = {imgproc::HSV2RGB_gpu_8u, 0, 0, 0, 0, imgproc::HSV2RGB_gpu_32f};\r
- static const func_t funcs_hls[] = {imgproc::HLS2RGB_gpu_8u, 0, 0, 0, 0, imgproc::HLS2RGB_gpu_32f};\r
-\r
- if (dcn <= 0) dcn = 3;\r
-\r
- CV_Assert((scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F));\r
-\r
- bidx = code == CV_HSV2BGR || code == CV_HLS2BGR ||\r
- code == CV_HSV2BGR_FULL || code == CV_HLS2BGR_FULL ? 0 : 2;\r
- int hrange = depth == CV_32F ? 360 : code == CV_HSV2BGR || code == CV_HSV2RGB ||\r
- code == CV_HLS2BGR || code == CV_HLS2RGB ? 180 : 255;\r
- \r
- dst.create(sz, CV_MAKETYPE(depth, dcn));\r
-\r
- if (code == CV_HSV2BGR || code == CV_HSV2RGB || code == CV_HSV2BGR_FULL || code == CV_HSV2RGB_FULL)\r
- funcs_hsv[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
- else\r
- funcs_hls[depth](src, scn, dst, dcn, bidx, hrange, stream);\r
- break;\r
- }\r
-\r
- default:\r
- CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
- }\r
- }\r
-}\r
-\r
-void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn)\r
-{\r
- cvtColor_caller(src, dst, code, dcn, 0);\r
-}\r
-\r
-void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, const Stream& stream)\r
-{\r
- cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream));\r
-}\r
-\r
-////////////////////////////////////////////////////////////////////////\r
// threshold\r
\r
double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh) \r
projects / profile / ivi / opencv.git / commitdiff