From 3ab2728da16e2a0440204603a6222ace4070a3bb Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 3 Aug 2011 12:10:36 +0000 Subject: [PATCH] gpu device layer code refactoring --- modules/gpu/CMakeLists.txt | 4 +- modules/gpu/src/color.cpp | 1728 +++++++++++++---- modules/gpu/src/cuda/brute_force_matcher.cu | 12 +- modules/gpu/src/cuda/color.cu | 1451 ++------------ modules/gpu/src/cuda/element_operations.cu | 242 +-- modules/gpu/src/cuda/filters.cu | 12 +- modules/gpu/src/cuda/hist.cu | 12 +- modules/gpu/src/cuda/match_template.cu | 12 +- modules/gpu/src/cuda/mathfunc.cu | 4 +- modules/gpu/src/cuda/matrix_reductions.cu | 116 +- modules/gpu/src/cuda/stereobp.cu | 6 +- modules/gpu/src/cuda/stereocsbp.cu | 20 +- modules/gpu/src/cuda/surf.cu | 50 +- .../opencv2/gpu/device/border_interpolate.hpp | 78 +- modules/gpu/src/opencv2/gpu/device/color.hpp | 221 +++ .../src/opencv2/gpu/device/datamov_utils.hpp | 68 +- .../src/opencv2/gpu/device/detail/color.hpp | 1037 ++++++++++ .../opencv2/gpu/device/detail/transform.hpp | 429 ++++ .../gpu/src/opencv2/gpu/device/functional.hpp | 338 ++++ .../gpu/device/{limits_gpu.hpp => limits.hpp} | 26 +- .../src/opencv2/gpu/device/saturate_cast.hpp | 220 +-- .../gpu/src/opencv2/gpu/device/transform.hpp | 421 +--- .../gpu/src/opencv2/gpu/device/utility.hpp | 206 ++ .../gpu/src/opencv2/gpu/device/vec_math.hpp | 287 +++ .../gpu/src/opencv2/gpu/device/vec_traits.hpp | 142 ++ .../gpu/src/opencv2/gpu/device/vecmath.hpp | 1097 ----------- modules/gpu/src/surf.cpp | 6 +- samples/gpu/performance/tests.cpp | 27 +- 28 files changed, 4506 insertions(+), 3766 deletions(-) create mode 100644 modules/gpu/src/opencv2/gpu/device/color.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/detail/color.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/detail/transform.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/functional.hpp rename modules/gpu/src/opencv2/gpu/device/{limits_gpu.hpp => limits.hpp} (92%) create mode 100644 modules/gpu/src/opencv2/gpu/device/utility.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/vec_math.hpp create mode 100644 modules/gpu/src/opencv2/gpu/device/vec_traits.hpp delete mode 100644 modules/gpu/src/opencv2/gpu/device/vecmath.hpp diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 84dfb856c3..eaa47a0f3f 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -23,7 +23,9 @@ source_group("Include" FILES ${lib_hdrs}) #file(GLOB lib_device_hdrs "include/opencv2/${name}/device/*.h*") file(GLOB lib_device_hdrs "src/opencv2/gpu/device/*.h*") +file(GLOB lib_device_hdrs_detail "src/opencv2/gpu/device/detail/*.h*") source_group("Device" FILES ${lib_device_hdrs}) +source_group("Device\\Detail" FILES ${lib_device_hdrs_detail}) if (HAVE_CUDA) file(GLOB_RECURSE ncv_srcs "src/nvidia/*.cpp") @@ -83,7 +85,7 @@ foreach(d ${DEPS}) endif() endforeach() -add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${lib_device_hdrs} ${ncv_srcs} ${ncv_hdrs} ${ncv_cuda} ${cuda_objs}) +add_library(${the_target} ${lib_srcs} ${lib_hdrs} ${lib_int_hdrs} ${lib_cuda} ${lib_cuda_hdrs} ${lib_device_hdrs} ${lib_device_hdrs_detail} ${ncv_srcs} ${ncv_hdrs} ${ncv_cuda} ${cuda_objs}) # For dynamic link numbering convenions set_target_properties(${the_target} PROPERTIES diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index 4d9bd50167..c539f3a6da 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -53,410 +53,1376 @@ void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu( 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); + #define OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name) \ + void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); + + #define OPENCV_GPU_DECLARE_CVTCOLOR_ALL(name) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _16u) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f) + + #define OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(name) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _8u) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _32f) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_8u) \ + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(name ## _full_32f) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_rgba) + + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr555) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr_to_bgr565) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr555) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgb_to_bgr565) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr555) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgra_to_bgr565) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr555) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(rgba_to_bgr565) + + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(gray_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr555) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(gray_to_bgr565) + + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr555_to_gray) + OPENCV_GPU_DECLARE_CVTCOLOR_ONE(bgr565_to_gray) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_gray) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_gray) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_gray) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_gray) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_yuv4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_yuv4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_yuv4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_yuv4) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(yuv4_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_YCrCb4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_YCrCb4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_YCrCb4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_YCrCb4) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(YCrCb4_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgb_to_xyz4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(rgba_to_xyz4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgr_to_xyz4) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(bgra_to_xyz4) + + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_ALL(xyz4_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hsv4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hsv4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hsv4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hsv4) + + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hsv4_to_bgra) + + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgb_to_hls4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(rgba_to_hls4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgr_to_hls4) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(bgra_to_hls4) + + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgb) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_rgba) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls_to_bgra) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgr) + OPENCV_GPU_DECLARE_CVTCOLOR_8U32F(hls4_to_bgra) + + #undef OPENCV_GPU_DECLARE_CVTCOLOR_ONE + #undef OPENCV_GPU_DECLARE_CVTCOLOR_ALL + #undef OPENCV_GPU_DECLARE_CVTCOLOR_8U32F +}}} - 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); +namespace +{ + typedef void (*gpu_func_t)(const DevMem2D& src, const DevMem2D& dst, 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 bgr_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgr_to_rgb_8u, 0, bgr_to_rgb_16u, 0, 0, bgr_to_rgb_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3); - 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); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); - 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); + funcs[src.depth()](src, dst, StreamAccessor::getStream(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 bgr_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgr_to_bgra_8u, 0, bgr_to_bgra_16u, 0, 0, bgr_to_bgra_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(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 bgr_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgr_to_rgba_8u, 0, bgr_to_rgba_16u, 0, 0, bgr_to_rgba_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3); - 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); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); - 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); + funcs[src.depth()](src, dst, StreamAccessor::getStream(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 bgra_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgra_to_bgr_8u, 0, bgra_to_bgr_16u, 0, 0, bgra_to_bgr_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 4); - 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); + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); - 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); -}}} + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } -namespace -{ - #undef R2Y - #undef G2Y - #undef B2Y + void bgra_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgra_to_rgb_8u, 0, bgra_to_rgb_16u, 0, 0, bgra_to_rgb_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgra_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgra_to_rgba_8u, 0, bgra_to_rgba_16u, 0, 0, bgra_to_rgba_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_8UC2); + + color::bgr_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_8UC2); + + color::bgr_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_8UC2); + + color::rgb_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_8UC2); + + color::rgb_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + } + + void bgra_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_8UC2); + + color::bgra_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + } + + void bgra_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_8UC2); + + color::bgra_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + } + + void rgba_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_8UC2); + + color::rgba_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + } + + void rgba_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_8UC2); + + color::rgba_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr555_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC3); + + color::bgr555_to_rgb(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr565_to_rgb(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC3); + + color::bgr565_to_rgb(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr555_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC3); + + color::bgr555_to_bgr(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr565_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC3); + + color::bgr565_to_bgr(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr555_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC4); + + color::bgr555_to_rgba(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr565_to_rgba(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC4); + + color::bgr565_to_rgba(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr555_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC4); + + color::bgr555_to_bgra(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr565_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC4); + + color::bgr565_to_bgra(src, dst, StreamAccessor::getStream(stream)); + } + + void gray_to_bgr(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {gray_to_bgr_8u, 0, gray_to_bgr_16u, 0, 0, gray_to_bgr_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 1); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 3)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void gray_to_bgra(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {gray_to_bgra_8u, 0, gray_to_bgra_16u, 0, 0, gray_to_bgra_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 1); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 4)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void gray_to_bgr555(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 1); + + dst.create(src.size(), CV_8UC2); + + color::gray_to_bgr555(src, dst, StreamAccessor::getStream(stream)); + } + + void gray_to_bgr565(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 1); + + dst.create(src.size(), CV_8UC2); + + color::gray_to_bgr565(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr555_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC1); + + color::bgr555_to_gray(src, dst, StreamAccessor::getStream(stream)); + } + + void bgr565_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 2); + + dst.create(src.size(), CV_8UC1); + + color::bgr565_to_gray(src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {rgb_to_gray_8u, 0, rgb_to_gray_16u, 0, 0, rgb_to_gray_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgr_to_gray_8u, 0, bgr_to_gray_16u, 0, 0, bgr_to_gray_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgba_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {rgba_to_gray_8u, 0, rgba_to_gray_16u, 0, 0, rgba_to_gray_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgra_to_gray(const GpuMat& src, GpuMat& dst, int, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[] = {bgra_to_gray_8u, 0, bgra_to_gray_16u, 0, 0, bgra_to_gray_32f}; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), 1)); + + funcs[src.depth()](src, dst, StreamAccessor::getStream(stream)); + } - enum - { - yuv_shift = 14, - xyz_shift = 12, - R2Y = 4899, - G2Y = 9617, - B2Y = 1868, - BLOCK_SIZE = 256 - }; -} + void rgb_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_yuv_8u, 0, rgb_to_yuv_16u, 0, 0, rgb_to_yuv_32f}, + {rgba_to_yuv_8u, 0, rgba_to_yuv_16u, 0, 0, rgba_to_yuv_32f} + }, + { + {rgb_to_yuv4_8u, 0, rgb_to_yuv4_16u, 0, 0, rgb_to_yuv4_32f}, + {rgba_to_yuv4_8u, 0, rgba_to_yuv4_16u, 0, 0, rgba_to_yuv4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); -namespace -{ - void cvtColor_caller(const GpuMat& src, GpuMat& dst, int code, int dcn, const cudaStream_t& stream) + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_yuv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_yuv_8u, 0, bgr_to_yuv_16u, 0, 0, bgr_to_yuv_32f}, + {bgra_to_yuv_8u, 0, bgra_to_yuv_16u, 0, 0, bgra_to_yuv_32f} + }, + { + {bgr_to_yuv4_8u, 0, bgr_to_yuv4_16u, 0, 0, bgr_to_yuv4_32f}, + {bgra_to_yuv4_8u, 0, bgra_to_yuv4_16u, 0, 0, bgra_to_yuv4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void yuv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {yuv_to_rgb_8u, 0, yuv_to_rgb_16u, 0, 0, yuv_to_rgb_32f}, + {yuv4_to_rgb_8u, 0, yuv4_to_rgb_16u, 0, 0, yuv4_to_rgb_32f} + }, + { + {yuv_to_rgba_8u, 0, yuv_to_rgba_16u, 0, 0, yuv_to_rgba_32f}, + {yuv4_to_rgba_8u, 0, yuv4_to_rgba_16u, 0, 0, yuv4_to_rgba_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void yuv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {yuv_to_bgr_8u, 0, yuv_to_bgr_16u, 0, 0, yuv_to_bgr_32f}, + {yuv4_to_bgr_8u, 0, yuv4_to_bgr_16u, 0, 0, yuv4_to_bgr_32f} + }, + { + {yuv_to_bgra_8u, 0, yuv_to_bgra_16u, 0, 0, yuv_to_bgra_32f}, + {yuv4_to_bgra_8u, 0, yuv4_to_bgra_16u, 0, 0, yuv4_to_bgra_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_YCrCb_8u, 0, rgb_to_YCrCb_16u, 0, 0, rgb_to_YCrCb_32f}, + {rgba_to_YCrCb_8u, 0, rgba_to_YCrCb_16u, 0, 0, rgba_to_YCrCb_32f} + }, + { + {rgb_to_YCrCb4_8u, 0, rgb_to_YCrCb4_16u, 0, 0, rgb_to_YCrCb4_32f}, + {rgba_to_YCrCb4_8u, 0, rgba_to_YCrCb4_16u, 0, 0, rgba_to_YCrCb4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_YCrCb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_YCrCb_8u, 0, bgr_to_YCrCb_16u, 0, 0, bgr_to_YCrCb_32f}, + {bgra_to_YCrCb_8u, 0, bgra_to_YCrCb_16u, 0, 0, bgra_to_YCrCb_32f} + }, + { + {bgr_to_YCrCb4_8u, 0, bgr_to_YCrCb4_16u, 0, 0, bgr_to_YCrCb4_32f}, + {bgra_to_YCrCb4_8u, 0, bgra_to_YCrCb4_16u, 0, 0, bgra_to_YCrCb4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void YCrCb_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {YCrCb_to_rgb_8u, 0, YCrCb_to_rgb_16u, 0, 0, YCrCb_to_rgb_32f}, + {YCrCb4_to_rgb_8u, 0, YCrCb4_to_rgb_16u, 0, 0, YCrCb4_to_rgb_32f} + }, + { + {YCrCb_to_rgba_8u, 0, YCrCb_to_rgba_16u, 0, 0, YCrCb_to_rgba_32f}, + {YCrCb4_to_rgba_8u, 0, YCrCb4_to_rgba_16u, 0, 0, YCrCb4_to_rgba_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void YCrCb_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {YCrCb_to_bgr_8u, 0, YCrCb_to_bgr_16u, 0, 0, YCrCb_to_bgr_32f}, + {YCrCb4_to_bgr_8u, 0, YCrCb4_to_bgr_16u, 0, 0, YCrCb4_to_bgr_32f} + }, + { + {YCrCb_to_bgra_8u, 0, YCrCb_to_bgra_16u, 0, 0, YCrCb_to_bgra_32f}, + {YCrCb4_to_bgra_8u, 0, YCrCb4_to_bgra_16u, 0, 0, YCrCb4_to_bgra_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_xyz_8u, 0, rgb_to_xyz_16u, 0, 0, rgb_to_xyz_32f}, + {rgba_to_xyz_8u, 0, rgba_to_xyz_16u, 0, 0, rgba_to_xyz_32f} + }, + { + {rgb_to_xyz4_8u, 0, rgb_to_xyz4_16u, 0, 0, rgb_to_xyz4_32f}, + {rgba_to_xyz4_8u, 0, rgba_to_xyz4_16u, 0, 0, rgba_to_xyz4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_xyz(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_xyz_8u, 0, bgr_to_xyz_16u, 0, 0, bgr_to_xyz_32f}, + {bgra_to_xyz_8u, 0, bgra_to_xyz_16u, 0, 0, bgra_to_xyz_32f} + }, + { + {bgr_to_xyz4_8u, 0, bgr_to_xyz4_16u, 0, 0, bgr_to_xyz4_32f}, + {bgra_to_xyz4_8u, 0, bgra_to_xyz4_16u, 0, 0, bgra_to_xyz4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void xyz_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {xyz_to_rgb_8u, 0, xyz_to_rgb_16u, 0, 0, xyz_to_rgb_32f}, + {xyz4_to_rgb_8u, 0, xyz4_to_rgb_16u, 0, 0, xyz4_to_rgb_32f} + }, + { + {xyz_to_rgba_8u, 0, xyz_to_rgba_16u, 0, 0, xyz_to_rgba_32f}, + {xyz4_to_rgba_8u, 0, xyz4_to_rgba_16u, 0, 0, xyz4_to_rgba_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void xyz_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {xyz_to_bgr_8u, 0, xyz_to_bgr_16u, 0, 0, xyz_to_bgr_32f}, + {xyz4_to_bgr_8u, 0, xyz4_to_bgr_16u, 0, 0, xyz4_to_bgr_32f} + }, + { + {xyz_to_bgra_8u, 0, xyz_to_bgra_16u, 0, 0, xyz_to_bgra_32f}, + {xyz4_to_bgra_8u, 0, xyz4_to_bgra_16u, 0, 0, xyz4_to_bgra_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_16U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_hsv_8u, 0, 0, 0, 0, rgb_to_hsv_32f}, + {rgba_to_hsv_8u, 0, 0, 0, 0, rgba_to_hsv_32f}, + }, + { + {rgb_to_hsv4_8u, 0, 0, 0, 0, rgb_to_hsv4_32f}, + {rgba_to_hsv4_8u, 0, 0, 0, 0, rgba_to_hsv4_32f}, + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_hsv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_hsv_8u, 0, 0, 0, 0, bgr_to_hsv_32f}, + {bgra_to_hsv_8u, 0, 0, 0, 0, bgra_to_hsv_32f} + }, + { + {bgr_to_hsv4_8u, 0, 0, 0, 0, bgr_to_hsv4_32f}, + {bgra_to_hsv4_8u, 0, 0, 0, 0, bgra_to_hsv4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hsv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hsv_to_rgb_8u, 0, 0, 0, 0, hsv_to_rgb_32f}, + {hsv4_to_rgb_8u, 0, 0, 0, 0, hsv4_to_rgb_32f} + }, + { + {hsv_to_rgba_8u, 0, 0, 0, 0, hsv_to_rgba_32f}, + {hsv4_to_rgba_8u, 0, 0, 0, 0, hsv4_to_rgba_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hsv_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hsv_to_bgr_8u, 0, 0, 0, 0, hsv_to_bgr_32f}, + {hsv4_to_bgr_8u, 0, 0, 0, 0, hsv4_to_bgr_32f} + }, + { + {hsv_to_bgra_8u, 0, 0, 0, 0, hsv_to_bgra_32f}, + {hsv4_to_bgra_8u, 0, 0, 0, 0, hsv4_to_bgra_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_hls_8u, 0, 0, 0, 0, rgb_to_hls_32f}, + {rgba_to_hls_8u, 0, 0, 0, 0, rgba_to_hls_32f}, + }, + { + {rgb_to_hls4_8u, 0, 0, 0, 0, rgb_to_hls4_32f}, + {rgba_to_hls4_8u, 0, 0, 0, 0, rgba_to_hls4_32f}, + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_hls(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_hls_8u, 0, 0, 0, 0, bgr_to_hls_32f}, + {bgra_to_hls_8u, 0, 0, 0, 0, bgra_to_hls_32f} + }, + { + {bgr_to_hls4_8u, 0, 0, 0, 0, bgr_to_hls4_32f}, + {bgra_to_hls4_8u, 0, 0, 0, 0, bgra_to_hls4_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hls_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hls_to_rgb_8u, 0, 0, 0, 0, hls_to_rgb_32f}, + {hls4_to_rgb_8u, 0, 0, 0, 0, hls4_to_rgb_32f} + }, + { + {hls_to_rgba_8u, 0, 0, 0, 0, hls_to_rgba_32f}, + {hls4_to_rgba_8u, 0, 0, 0, 0, hls4_to_rgba_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hls_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hls_to_bgr_8u, 0, 0, 0, 0, hls_to_bgr_32f}, + {hls4_to_bgr_8u, 0, 0, 0, 0, hls4_to_bgr_32f} + }, + { + {hls_to_bgra_8u, 0, 0, 0, 0, hls_to_bgra_32f}, + {hls4_to_bgra_8u, 0, 0, 0, 0, hls4_to_bgra_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_hsv_full_8u, 0, 0, 0, 0, rgb_to_hsv_full_32f}, + {rgba_to_hsv_full_8u, 0, 0, 0, 0, rgba_to_hsv_full_32f}, + }, + { + {rgb_to_hsv4_full_8u, 0, 0, 0, 0, rgb_to_hsv4_full_32f}, + {rgba_to_hsv4_full_8u, 0, 0, 0, 0, rgba_to_hsv4_full_32f}, + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_hsv_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {bgr_to_hsv_full_8u, 0, 0, 0, 0, bgr_to_hsv_full_32f}, + {bgra_to_hsv_full_8u, 0, 0, 0, 0, bgra_to_hsv_full_32f} + }, + { + {bgr_to_hsv4_full_8u, 0, 0, 0, 0, bgr_to_hsv4_full_32f}, + {bgra_to_hsv4_full_8u, 0, 0, 0, 0, bgra_to_hsv4_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hsv_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hsv_to_rgb_full_8u, 0, 0, 0, 0, hsv_to_rgb_full_32f}, + {hsv4_to_rgb_full_8u, 0, 0, 0, 0, hsv4_to_rgb_full_32f} + }, + { + {hsv_to_rgba_full_8u, 0, 0, 0, 0, hsv_to_rgba_full_32f}, + {hsv4_to_rgba_full_8u, 0, 0, 0, 0, hsv4_to_rgba_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hsv_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hsv_to_bgr_full_8u, 0, 0, 0, 0, hsv_to_bgr_full_32f}, + {hsv4_to_bgr_full_8u, 0, 0, 0, 0, hsv4_to_bgr_full_32f} + }, + { + {hsv_to_bgra_full_8u, 0, 0, 0, 0, hsv_to_bgra_full_32f}, + {hsv4_to_bgra_full_8u, 0, 0, 0, 0, hsv4_to_bgra_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void rgb_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) { - Size sz = src.size(); - int scn = src.channels(), depth = src.depth(), bidx; + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {rgb_to_hls_full_8u, 0, 0, 0, 0, rgb_to_hls_full_32f}, + {rgba_to_hls_full_8u, 0, 0, 0, 0, rgba_to_hls_full_32f}, + }, + { + {rgb_to_hls4_full_8u, 0, 0, 0, 0, rgb_to_hls4_full_32f}, + {rgba_to_hls4_full_8u, 0, 0, 0, 0, rgba_to_hls4_full_32f}, + } + }; + + if (dcn <= 0) dcn = 3; - CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32F); + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); - switch (code) + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void bgr_to_hls_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = { - 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, 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 : 256; - - 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" ); - } + { + {bgr_to_hls_full_8u, 0, 0, 0, 0, bgr_to_hls_full_32f}, + {bgra_to_hls_full_8u, 0, 0, 0, 0, bgra_to_hls_full_32f} + }, + { + {bgr_to_hls4_full_8u, 0, 0, 0, 0, bgr_to_hls4_full_32f}, + {bgra_to_hls4_full_8u, 0, 0, 0, 0, bgra_to_hls4_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hls_to_rgb_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hls_to_rgb_full_8u, 0, 0, 0, 0, hls_to_rgb_full_32f}, + {hls4_to_rgb_full_8u, 0, 0, 0, 0, hls4_to_rgb_full_32f} + }, + { + {hls_to_rgba_full_8u, 0, 0, 0, 0, hls_to_rgba_full_32f}, + {hls4_to_rgba_full_8u, 0, 0, 0, 0, hls4_to_rgba_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); + } + + void hls_to_bgr_full(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + using namespace cv::gpu::color; + static const gpu_func_t funcs[2][2][6] = + { + { + {hls_to_bgr_full_8u, 0, 0, 0, 0, hls_to_bgr_full_32f}, + {hls4_to_bgr_full_8u, 0, 0, 0, 0, hls4_to_bgr_full_32f} + }, + { + {hls_to_bgra_full_8u, 0, 0, 0, 0, hls_to_bgra_full_32f}, + {hls4_to_bgra_full_8u, 0, 0, 0, 0, hls4_to_bgra_full_32f} + } + }; + + if (dcn <= 0) dcn = 3; + + CV_Assert(src.depth() == CV_8U || src.depth() == CV_32F); + CV_Assert(src.channels() == 3 || src.channels() == 4); + CV_Assert(dcn == 3 || dcn == 4); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } } void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream) { - cvtColor_caller(src, dst, code, dcn, StreamAccessor::getStream(stream)); + typedef void (*func_t)(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream); + static const func_t funcs[] = + { + bgr_to_bgra, // CV_BGR2BGRA =0 + bgra_to_bgr, // CV_BGRA2BGR =1 + bgr_to_rgba, // CV_BGR2RGBA =2 + bgra_to_rgb, // CV_RGBA2BGR =3 + bgr_to_rgb, // CV_BGR2RGB =4 + bgra_to_rgba, // CV_BGRA2RGBA =5 + + bgr_to_gray, // CV_BGR2GRAY =6 + rgb_to_gray, // CV_RGB2GRAY =7 + gray_to_bgr, // CV_GRAY2BGR =8 + gray_to_bgra, // CV_GRAY2BGRA =9 + bgra_to_gray, // CV_BGRA2GRAY =10 + rgba_to_gray, // CV_RGBA2GRAY =11 + + bgr_to_bgr565, // CV_BGR2BGR565 =12 + rgb_to_bgr565, // CV_RGB2BGR565 =13 + bgr565_to_bgr, // CV_BGR5652BGR =14 + bgr565_to_rgb, // CV_BGR5652RGB =15 + bgra_to_bgr565, // CV_BGRA2BGR565 =16 + rgba_to_bgr565, // CV_RGBA2BGR565 =17 + bgr565_to_bgra, // CV_BGR5652BGRA =18 + bgr565_to_rgba, // CV_BGR5652RGBA =19 + + gray_to_bgr565, // CV_GRAY2BGR565 =20 + bgr565_to_gray, // CV_BGR5652GRAY =21 + + bgr_to_bgr555, // CV_BGR2BGR555 =22 + rgb_to_bgr555, // CV_RGB2BGR555 =23 + bgr555_to_bgr, // CV_BGR5552BGR =24 + bgr555_to_rgb, // CV_BGR5552RGB =25 + bgra_to_bgr555, // CV_BGRA2BGR555 =26 + rgba_to_bgr555, // CV_RGBA2BGR555 =27 + bgr555_to_bgra, // CV_BGR5552BGRA =28 + bgr555_to_rgba, // CV_BGR5552RGBA =29 + + gray_to_bgr555, // CV_GRAY2BGR555 =30 + bgr555_to_gray, // CV_BGR5552GRAY =31 + + bgr_to_xyz, // CV_BGR2XYZ =32 + rgb_to_xyz, // CV_RGB2XYZ =33 + xyz_to_bgr, // CV_XYZ2BGR =34 + xyz_to_rgb, // CV_XYZ2RGB =35 + + bgr_to_YCrCb, // CV_BGR2YCrCb =36 + rgb_to_YCrCb, // CV_RGB2YCrCb =37 + YCrCb_to_bgr, // CV_YCrCb2BGR =38 + YCrCb_to_rgb, // CV_YCrCb2RGB =39 + + bgr_to_hsv, // CV_BGR2HSV =40 + rgb_to_hsv, // CV_RGB2HSV =41 + + 0, // =42 + 0, // =43 + + 0, // CV_BGR2Lab =44 + 0, // CV_RGB2Lab =45 + + 0, // CV_BayerBG2BGR =46 + 0, // CV_BayerGB2BGR =47 + 0, // CV_BayerRG2BGR =48 + 0, // CV_BayerGR2BGR =49 + + 0, // CV_BGR2Luv =50 + 0, // CV_RGB2Luv =51 + + bgr_to_hls, // CV_BGR2HLS =52 + rgb_to_hls, // CV_RGB2HLS =53 + + hsv_to_bgr, // CV_HSV2BGR =54 + bgr_to_rgb, // CV_HSV2RGB =55 + + 0, // CV_Lab2BGR =56 + 0, // CV_Lab2RGB =57 + 0, // CV_Luv2BGR =58 + 0, // CV_Luv2RGB =59 + + hls_to_bgr, // CV_HLS2BGR =60 + hls_to_rgb, // CV_HLS2RGB =61 + + 0, // CV_BayerBG2BGR_VNG =62 + 0, // CV_BayerGB2BGR_VNG =63 + 0, // CV_BayerRG2BGR_VNG =64 + 0, // CV_BayerGR2BGR_VNG =65 + + bgr_to_hsv_full, // CV_BGR2HSV_FULL = 66 + rgb_to_hsv_full, // CV_RGB2HSV_FULL = 67 + bgr_to_hls_full, // CV_BGR2HLS_FULL = 68 + rgb_to_hls_full, // CV_RGB2HLS_FULL = 69 + + hsv_to_bgr_full, // CV_HSV2BGR_FULL = 70 + hsv_to_rgb_full, // CV_HSV2RGB_FULL = 71 + hls_to_bgr_full, // CV_HLS2BGR_FULL = 72 + hls_to_rgb_full, // CV_HLS2RGB_FULL = 73 + + 0, // CV_LBGR2Lab = 74 + 0, // CV_LRGB2Lab = 75 + 0, // CV_LBGR2Luv = 76 + 0, // CV_LRGB2Luv = 77 + + 0, // CV_Lab2LBGR = 78 + 0, // CV_Lab2LRGB = 79 + 0, // CV_Luv2LBGR = 80 + 0, // CV_Luv2LRGB = 81 + + bgr_to_yuv, // CV_BGR2YUV = 82 + rgb_to_yuv, // CV_RGB2YUV = 83 + yuv_to_bgr, // CV_YUV2BGR = 84 + yuv_to_rgb, // CV_YUV2RGB = 85 + + 0, // CV_BayerBG2GRAY = 86 + 0, // CV_BayerGB2GRAY = 87 + 0, // CV_BayerRG2GRAY = 88 + 0, // CV_BayerGR2GRAY = 89 + + 0, // CV_YUV420i2RGB = 90 + 0, // CV_YUV420i2BGR = 91 + 0, // CV_YUV420sp2RGB = 92 + 0 // CV_YUV420sp2BGR = 93 + }; + + CV_Assert(code < 94); + + func_t func = funcs[code]; + + if (func == 0) + CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); + + func(src, dst, dcn, stream); } #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/cuda/brute_force_matcher.cu b/modules/gpu/src/cuda/brute_force_matcher.cu index 27b8e530a4..17c5c802f7 100644 --- a/modules/gpu/src/cuda/brute_force_matcher.cu +++ b/modules/gpu/src/cuda/brute_force_matcher.cu @@ -41,7 +41,7 @@ //M*/ #include "internal_shared.hpp" -#include "opencv2/gpu/device/limits_gpu.hpp" +#include "opencv2/gpu/device/limits.hpp" #include "opencv2/gpu/device/datamov_utils.hpp" using namespace cv::gpu; @@ -565,7 +565,7 @@ namespace cv { namespace gpu { namespace bfmatcher int myBestTrainIdx = -1; int myBestImgIdx = -1; - typename Dist::ResultType myMin = numeric_limits_gpu::max(); + typename Dist::ResultType myMin = numeric_limits::max(); { typename Dist::ResultType* sdiff_row = smem + BLOCK_DIM_X * threadIdx.y; @@ -821,7 +821,7 @@ namespace cv { namespace gpu { namespace bfmatcher { const T* trainDescs = trainDescs_.ptr(trainIdx); - typename Dist::ResultType myDist = numeric_limits_gpu::max(); + typename Dist::ResultType myDist = numeric_limits::max(); if (mask(queryIdx, trainIdx)) { @@ -932,7 +932,7 @@ namespace cv { namespace gpu { namespace bfmatcher { const int tid = threadIdx.x; - T myMin = numeric_limits_gpu::max(); + T myMin = numeric_limits::max(); int myMinIdx = -1; for (int i = tid; i < n; i += BLOCK_SIZE) @@ -1007,10 +1007,10 @@ namespace cv { namespace gpu { namespace bfmatcher if (threadIdx.x == 0) { float dist = sdist[0]; - if (dist < numeric_limits_gpu::max()) + if (dist < numeric_limits::max()) { int bestIdx = strainIdx[0]; - allDist[bestIdx] = numeric_limits_gpu::max(); + allDist[bestIdx] = numeric_limits::max(); trainIdx[i] = bestIdx; distance[i] = dist; } diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index 9fcd520619..2755fffba8 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -41,1305 +41,168 @@ //M*/ #include "internal_shared.hpp" -#include "opencv2/gpu/device/saturate_cast.hpp" -#include "opencv2/gpu/device/vecmath.hpp" -#include "opencv2/gpu/device/limits_gpu.hpp" #include "opencv2/gpu/device/transform.hpp" +#include "opencv2/gpu/device/color.hpp" using namespace cv::gpu; using namespace cv::gpu::device; -#ifndef CV_DESCALE -#define CV_DESCALE(x, n) (((x) + (1 << ((n)-1))) >> (n)) -#endif - namespace cv { namespace gpu { namespace color { - template struct ColorChannel - { - typedef float worktype_f; - static __device__ __forceinline__ T max() { return numeric_limits_gpu::max(); } - static __device__ __forceinline__ T half() { return (T)(max()/2 + 1); } - }; - template<> struct ColorChannel - { - typedef float worktype_f; - static __device__ __forceinline__ float max() { return 1.f; } - static __device__ __forceinline__ float half() { return 0.5f; } - }; - - template - __device__ __forceinline__ void setAlpha(typename TypeVec::vec_t& vec, T val) - { - } - template - __device__ __forceinline__ void setAlpha(typename TypeVec::vec_t& vec, T val) - { - vec.w = val; - } - template - __device__ __forceinline__ T getAlpha(const typename TypeVec::vec_t& vec) - { - return ColorChannel::max(); - } - template - __device__ __forceinline__ T getAlpha(const typename TypeVec::vec_t& vec) - { - return vec.w; - } - - template - void callConvert(const DevMem2D& src, const DevMem2D& dst, const Cvt& cvt, cudaStream_t stream) - { - typedef typename Cvt::src_t src_t; - typedef typename Cvt::dst_t dst_t; - - transform((DevMem2D_)src, (DevMem2D_)dst, cvt, stream); - } - -////////////////// Various 3/4-channel to 3/4-channel RGB transformations ///////////////// - - template - struct RGB2RGB - { - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - explicit RGB2RGB(int bidx) : bidx(bidx) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - - dst.x = (&src.x)[bidx]; - dst.y = src.y; - dst.z = (&src.x)[bidx ^ 2]; - setAlpha(dst, getAlpha(src)); - - return dst; - } - - private: - int bidx; - }; - - template - void RGB2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - RGB2RGB cvt(bidx); - callConvert(src, dst, cvt, stream); - } - - void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) - { - typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = - { - {RGB2RGB_caller, RGB2RGB_caller}, - {RGB2RGB_caller, RGB2RGB_caller} - }; - - RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); - } - - void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) - { - typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = - { - {RGB2RGB_caller, RGB2RGB_caller}, - {RGB2RGB_caller, RGB2RGB_caller} - }; - - RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); - } - - void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) - { - typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = - { - {RGB2RGB_caller, RGB2RGB_caller}, - {RGB2RGB_caller, RGB2RGB_caller} - }; - - RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); - } - -/////////// Transforming 16-bit (565 or 555) RGB to/from 24/32-bit (888[8]) RGB ////////// - - template struct RGB5x52RGBConverter; - template <> struct RGB5x52RGBConverter<5> - { - template - static __device__ __forceinline__ void cvt(uint src, D& dst, int bidx) - { - (&dst.x)[bidx] = (uchar)(src << 3); - dst.y = (uchar)((src >> 2) & ~7); - (&dst.x)[bidx ^ 2] = (uchar)((src >> 7) & ~7); - setAlpha(dst, (uchar)(src & 0x8000 ? 255 : 0)); - } - }; - template <> struct RGB5x52RGBConverter<6> - { - template - static __device__ __forceinline__ void cvt(uint src, D& dst, int bidx) - { - (&dst.x)[bidx] = (uchar)(src << 3); - dst.y = (uchar)((src >> 3) & ~3); - (&dst.x)[bidx ^ 2] = (uchar)((src >> 8) & ~7); - setAlpha(dst, (uchar)(255)); - } - }; - - template struct RGB5x52RGB - { - typedef ushort src_t; - typedef typename TypeVec::vec_t dst_t; - - explicit RGB5x52RGB(int bidx) : bidx(bidx) {} - - __device__ __forceinline__ dst_t operator()(ushort src) const - { - dst_t dst; - RGB5x52RGBConverter::cvt((uint)src, dst, bidx); - return dst; - } - - private: - int bidx; - }; - - template struct RGB2RGB5x5Converter; - template<> struct RGB2RGB5x5Converter<6> - { - template - static __device__ __forceinline__ ushort cvt(const T& src, int bidx) - { - return (ushort)(((&src.x)[bidx] >> 3) | ((src.y & ~3) << 3) | (((&src.x)[bidx^2] & ~7) << 8)); - } - }; - template<> struct RGB2RGB5x5Converter<5> - { - static __device__ __forceinline__ ushort cvt(const uchar3& src, int bidx) - { - return (ushort)(((&src.x)[bidx] >> 3) | ((src.y & ~7) << 2) | (((&src.x)[bidx^2] & ~7) << 7)); - } - static __device__ __forceinline__ ushort cvt(const uchar4& src, int bidx) - { - return (ushort)(((&src.x)[bidx] >> 3) | ((src.y & ~7) << 2) | (((&src.x)[bidx^2] & ~7) << 7) | (src.w ? 0x8000 : 0)); - } - }; - - template struct RGB2RGB5x5 - { - typedef typename TypeVec::vec_t src_t; - typedef ushort dst_t; - - explicit RGB2RGB5x5(int bidx) : bidx(bidx) {} - - __device__ __forceinline__ ushort operator()(const src_t& src) - { - return RGB2RGB5x5Converter::cvt(src, bidx); - } - - private: - int bidx; - }; - - template - void RGB5x52RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - RGB5x52RGB cvt(bidx); - callConvert(src, dst, cvt, stream); - } - - void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) - { - typedef void (*RGB5x52RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - static const RGB5x52RGB_caller_t RGB5x52RGB_callers[2][2] = - { - {RGB5x52RGB_caller<5, 3>, RGB5x52RGB_caller<5, 4>}, - {RGB5x52RGB_caller<6, 3>, RGB5x52RGB_caller<6, 4>} - }; - - RGB5x52RGB_callers[green_bits - 5][dstcn - 3](src, dst, bidx, stream); - } - - template - void RGB2RGB5x5_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - RGB2RGB5x5 cvt(bidx); - callConvert(src, dst, cvt, stream); - } - - void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream) - { - typedef void (*RGB2RGB5x5_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - static const RGB2RGB5x5_caller_t RGB2RGB5x5_callers[2][2] = - { - {RGB2RGB5x5_caller<3, 5>, RGB2RGB5x5_caller<3, 6>}, - {RGB2RGB5x5_caller<4, 5>, RGB2RGB5x5_caller<4, 6>} - }; - - RGB2RGB5x5_callers[srccn - 3][green_bits - 5](src, dst, bidx, stream); - } - -///////////////////////////////// Grayscale to Color //////////////////////////////// - - template struct Gray2RGB - { - typedef T src_t; - typedef typename TypeVec::vec_t dst_t; - - __device__ __forceinline__ dst_t operator()(const T& src) const - { - dst_t dst; - - dst.z = dst.y = dst.x = src; - setAlpha(dst, ColorChannel::max()); - - return dst; - } - }; - - template struct Gray2RGB5x5Converter; - template<> struct Gray2RGB5x5Converter<6> - { - static __device__ __forceinline__ ushort cvt(uint t) - { - return (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); - } - }; - template<> struct Gray2RGB5x5Converter<5> - { - static __device__ __forceinline__ ushort cvt(uint t) - { - t >>= 3; - return (ushort)(t | (t << 5) | (t << 10)); - } - }; - - template struct Gray2RGB5x5 - { - typedef uchar src_t; - typedef ushort dst_t; - - __device__ __forceinline__ ushort operator()(uchar src) const - { - return Gray2RGB5x5Converter::cvt((uint)src); - } - }; - - template - void Gray2RGB_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) - { - Gray2RGB cvt; - callConvert(src, dst, cvt, stream); - } - - void Gray2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream) - { - typedef void (*Gray2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); - static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; - - Gray2RGB_callers[dstcn - 3](src, dst, stream); - } - - void Gray2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream) - { - typedef void (*Gray2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); - static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; - - Gray2RGB_callers[dstcn - 3](src, dst, stream); - } - - void Gray2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream) - { - typedef void (*Gray2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); - static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; - - Gray2RGB_callers[dstcn - 3](src, dst, stream); - } - - template - void Gray2RGB5x5_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) - { - Gray2RGB5x5 cvt; - callConvert(src, dst, cvt, stream); - } - - void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream) - { - typedef void (*Gray2RGB5x5_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); - static const Gray2RGB5x5_caller_t Gray2RGB5x5_callers[2] = - { - Gray2RGB5x5_caller<5>, Gray2RGB5x5_caller<6> - }; - - Gray2RGB5x5_callers[green_bits - 5](src, dst, stream); - } - -///////////////////////////////// Color to Grayscale //////////////////////////////// - - #undef R2Y - #undef G2Y - #undef B2Y - - enum - { - yuv_shift = 14, - xyz_shift = 12, - R2Y = 4899, - G2Y = 9617, - B2Y = 1868, - BLOCK_SIZE = 256 - }; - - template struct RGB5x52GrayConverter; - template<> struct RGB5x52GrayConverter<6> - { - static __device__ __forceinline__ uchar cvt(uint t) - { - return (uchar)CV_DESCALE(((t << 3) & 0xf8) * B2Y + ((t >> 3) & 0xfc) * G2Y + ((t >> 8) & 0xf8) * R2Y, yuv_shift); - } - }; - template<> struct RGB5x52GrayConverter<5> - { - static __device__ __forceinline__ uchar cvt(uint t) - { - return (uchar)CV_DESCALE(((t << 3) & 0xf8) * B2Y + ((t >> 2) & 0xf8) * G2Y + ((t >> 7) & 0xf8) * R2Y, yuv_shift); - } - }; - - template struct RGB5x52Gray - { - typedef ushort src_t; - typedef uchar dst_t; - - __device__ __forceinline__ uchar operator()(ushort src) const - { - return RGB5x52GrayConverter::cvt((uint)src); - } - }; - - template - __device__ __forceinline__ T RGB2GrayConvert(const T* src, int bidx) - { - return (T)CV_DESCALE((unsigned)(src[bidx] * B2Y + src[1] * G2Y + src[bidx^2] * R2Y), yuv_shift); - } - __device__ __forceinline__ float RGB2GrayConvert(const float* src, int bidx) - { - const float cr = 0.299f; - const float cg = 0.587f; - const float cb = 0.114f; - - return src[bidx] * cb + src[1] * cg + src[bidx^2] * cr; - } - - template struct RGB2Gray - { - typedef typename TypeVec::vec_t src_t; - typedef T dst_t; - - explicit RGB2Gray(int bidx) : bidx(bidx) {} - - __device__ __forceinline__ T operator()(const src_t& src) - { - return RGB2GrayConvert(&src.x, bidx); - } - - private: - int bidx; - }; - - template - void RGB2Gray_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - RGB2Gray cvt(bidx); - callConvert(src, dst, cvt, stream); - } - - void RGB2Gray_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - typedef void (*RGB2Gray_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - RGB2Gray_caller_t RGB2Gray_callers[] = {RGB2Gray_caller, RGB2Gray_caller}; - - RGB2Gray_callers[srccn - 3](src, dst, bidx, stream); - } - - void RGB2Gray_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - typedef void (*RGB2Gray_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - RGB2Gray_caller_t RGB2Gray_callers[] = {RGB2Gray_caller, RGB2Gray_caller}; - - RGB2Gray_callers[srccn - 3](src, dst, bidx, stream); - } - - void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream) - { - typedef void (*RGB2Gray_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); - RGB2Gray_caller_t RGB2Gray_callers[] = {RGB2Gray_caller, RGB2Gray_caller}; - - RGB2Gray_callers[srccn - 3](src, dst, bidx, stream); - } - - template - void RGB5x52Gray_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) - { - RGB5x52Gray cvt; - callConvert(src, dst, cvt, stream); - } - - void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream) - { - typedef void (*RGB5x52Gray_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); - static const RGB5x52Gray_caller_t RGB5x52Gray_callers[2] = - { - RGB5x52Gray_caller<5>, RGB5x52Gray_caller<6> - }; - - RGB5x52Gray_callers[green_bits - 5](src, dst, stream); - } - -///////////////////////////////////// RGB <-> YCrCb ////////////////////////////////////// - - __constant__ int cYCrCbCoeffs_i[5]; - __constant__ float cYCrCbCoeffs_f[5]; - - template - __device__ __forceinline__ void RGB2YCrCbConvert(const T* src, D& dst, int bidx) - { - const int delta = ColorChannel::half() * (1 << yuv_shift); - - const int Y = CV_DESCALE(src[0] * cYCrCbCoeffs_i[0] + src[1] * cYCrCbCoeffs_i[1] + src[2] * cYCrCbCoeffs_i[2], yuv_shift); - const int Cr = CV_DESCALE((src[bidx^2] - Y) * cYCrCbCoeffs_i[3] + delta, yuv_shift); - const int Cb = CV_DESCALE((src[bidx] - Y) * cYCrCbCoeffs_i[4] + delta, yuv_shift); - - dst.x = saturate_cast(Y); - dst.y = saturate_cast(Cr); - dst.z = saturate_cast(Cb); - } - template - static __device__ __forceinline__ void RGB2YCrCbConvert(const float* src, D& dst, int bidx) - { - dst.x = src[0] * cYCrCbCoeffs_f[0] + src[1] * cYCrCbCoeffs_f[1] + src[2] * cYCrCbCoeffs_f[2]; - dst.y = (src[bidx^2] - dst.x) * cYCrCbCoeffs_f[3] + ColorChannel::half(); - dst.z = (src[bidx] - dst.x) * cYCrCbCoeffs_f[4] + ColorChannel::half(); - } - - template struct RGB2YCrCbBase - { - typedef int coeff_t; - - explicit RGB2YCrCbBase(const coeff_t coeffs[5]) - { - cudaSafeCall( cudaMemcpyToSymbol(cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) ); - } - }; - template<> struct RGB2YCrCbBase - { - typedef float coeff_t; - - explicit RGB2YCrCbBase(const coeff_t coeffs[5]) - { - cudaSafeCall( cudaMemcpyToSymbol(cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) ); - } - }; - template struct RGB2YCrCb : RGB2YCrCbBase - { - typedef typename RGB2YCrCbBase::coeff_t coeff_t; - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - RGB2YCrCb(int bidx, const coeff_t coeffs[5]) : RGB2YCrCbBase(coeffs), bidx(bidx) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - RGB2YCrCbConvert(&src.x, dst, bidx); - return dst; - } - - private: - int bidx; - }; - - template - __device__ __forceinline__ void YCrCb2RGBConvert(const T& src, D* dst, int bidx) - { - const int b = src.x + CV_DESCALE((src.z - ColorChannel::half()) * cYCrCbCoeffs_i[3], yuv_shift); - const int g = src.x + CV_DESCALE((src.z - ColorChannel::half()) * cYCrCbCoeffs_i[2] + (src.y - ColorChannel::half()) * cYCrCbCoeffs_i[1], yuv_shift); - const int r = src.x + CV_DESCALE((src.y - ColorChannel::half()) * cYCrCbCoeffs_i[0], yuv_shift); - - dst[bidx] = saturate_cast(b); - dst[1] = saturate_cast(g); - dst[bidx^2] = saturate_cast(r); - } - template - __device__ __forceinline__ void YCrCb2RGBConvert(const T& src, float* dst, int bidx) - { - dst[bidx] = src.x + (src.z - ColorChannel::half()) * cYCrCbCoeffs_f[3]; - dst[1] = src.x + (src.z - ColorChannel::half()) * cYCrCbCoeffs_f[2] + (src.y - ColorChannel::half()) * cYCrCbCoeffs_f[1]; - dst[bidx^2] = src.x + (src.y - ColorChannel::half()) * cYCrCbCoeffs_f[0]; - } - - template struct YCrCb2RGBBase - { - typedef int coeff_t; - - explicit YCrCb2RGBBase(const coeff_t coeffs[4]) - { - cudaSafeCall( cudaMemcpyToSymbol(cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) ); - } - }; - template<> struct YCrCb2RGBBase - { - typedef float coeff_t; - - explicit YCrCb2RGBBase(const coeff_t coeffs[4]) - { - cudaSafeCall( cudaMemcpyToSymbol(cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) ); - } - }; - template struct YCrCb2RGB : YCrCb2RGBBase - { - typedef typename YCrCb2RGBBase::coeff_t coeff_t; - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - YCrCb2RGB(int bidx, const coeff_t coeffs[4]) : YCrCb2RGBBase(coeffs), bidx(bidx) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - - YCrCb2RGBConvert(src, &dst.x, bidx); - setAlpha(dst, ColorChannel::max()); - - return dst; - } - - private: - int bidx; - }; - - template - void RGB2YCrCb_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef typename RGB2YCrCb::coeff_t coeff_t; - RGB2YCrCb cvt(bidx, (const coeff_t*)coeffs); - callConvert(src, dst, cvt, stream); - } - - void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = - { - {RGB2YCrCb_caller, RGB2YCrCb_caller}, - {RGB2YCrCb_caller, RGB2YCrCb_caller} - }; - - RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - - void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = - { - {RGB2YCrCb_caller, RGB2YCrCb_caller}, - {RGB2YCrCb_caller, RGB2YCrCb_caller} - }; - - RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - - void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = - { - {RGB2YCrCb_caller, RGB2YCrCb_caller}, - {RGB2YCrCb_caller, RGB2YCrCb_caller} - }; - - RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - - template - void YCrCb2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef typename YCrCb2RGB::coeff_t coeff_t; - YCrCb2RGB cvt(bidx, (const coeff_t*)coeffs); - callConvert(src, dst, cvt, stream); - } - - void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = - { - {YCrCb2RGB_caller, YCrCb2RGB_caller}, - {YCrCb2RGB_caller, YCrCb2RGB_caller} - }; - - YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - - void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = - { - {YCrCb2RGB_caller, YCrCb2RGB_caller}, - {YCrCb2RGB_caller, YCrCb2RGB_caller} - }; - - YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - - void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const void* coeffs, cudaStream_t stream) - { - typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, const void* coeffs, cudaStream_t stream); - static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = - { - {YCrCb2RGB_caller, YCrCb2RGB_caller}, - {YCrCb2RGB_caller, YCrCb2RGB_caller} - }; - - YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, coeffs, stream); - } - -////////////////////////////////////// RGB <-> XYZ /////////////////////////////////////// - - __constant__ int cXYZ_D65i[9]; - __constant__ float cXYZ_D65f[9]; - - template - __device__ __forceinline__ void RGB2XYZConvert(const T* src, D& dst) - { - dst.x = saturate_cast(CV_DESCALE(src[0] * cXYZ_D65i[0] + src[1] * cXYZ_D65i[1] + src[2] * cXYZ_D65i[2], xyz_shift)); - dst.y = saturate_cast(CV_DESCALE(src[0] * cXYZ_D65i[3] + src[1] * cXYZ_D65i[4] + src[2] * cXYZ_D65i[5], xyz_shift)); - dst.z = saturate_cast(CV_DESCALE(src[0] * cXYZ_D65i[6] + src[1] * cXYZ_D65i[7] + src[2] * cXYZ_D65i[8], xyz_shift)); - } - template - __device__ __forceinline__ void RGB2XYZConvert(const float* src, D& dst) - { - dst.x = src[0] * cXYZ_D65f[0] + src[1] * cXYZ_D65f[1] + src[2] * cXYZ_D65f[2]; - dst.y = src[0] * cXYZ_D65f[3] + src[1] * cXYZ_D65f[4] + src[2] * cXYZ_D65f[5]; - dst.z = src[0] * cXYZ_D65f[6] + src[1] * cXYZ_D65f[7] + src[2] * cXYZ_D65f[8]; - } - - template struct RGB2XYZBase - { - typedef int coeff_t; - - explicit RGB2XYZBase(const coeff_t coeffs[9]) - { - cudaSafeCall( cudaMemcpyToSymbol(cXYZ_D65i, coeffs, 9 * sizeof(int)) ); - } - }; - template <> struct RGB2XYZBase - { - typedef float coeff_t; - - explicit RGB2XYZBase(const coeff_t coeffs[9]) - { - cudaSafeCall( cudaMemcpyToSymbol(cXYZ_D65f, coeffs, 9 * sizeof(float)) ); - } - }; - template struct RGB2XYZ : RGB2XYZBase - { - typedef typename RGB2XYZBase::coeff_t coeff_t; - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - explicit RGB2XYZ(const coeff_t coeffs[9]) : RGB2XYZBase(coeffs) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - RGB2XYZConvert(&src.x, dst); - return dst; - } - }; - - template - __device__ __forceinline__ void XYZ2RGBConvert(const T& src, D* dst) - { - dst[0] = saturate_cast(CV_DESCALE(src.x * cXYZ_D65i[0] + src.y * cXYZ_D65i[1] + src.z * cXYZ_D65i[2], xyz_shift)); - dst[1] = saturate_cast(CV_DESCALE(src.x * cXYZ_D65i[3] + src.y * cXYZ_D65i[4] + src.z * cXYZ_D65i[5], xyz_shift)); - dst[2] = saturate_cast(CV_DESCALE(src.x * cXYZ_D65i[6] + src.y * cXYZ_D65i[7] + src.z * cXYZ_D65i[8], xyz_shift)); - } - template - __device__ __forceinline__ void XYZ2RGBConvert(const T& src, float* dst) - { - dst[0] = src.x * cXYZ_D65f[0] + src.y * cXYZ_D65f[1] + src.z * cXYZ_D65f[2]; - dst[1] = src.x * cXYZ_D65f[3] + src.y * cXYZ_D65f[4] + src.z * cXYZ_D65f[5]; - dst[2] = src.x * cXYZ_D65f[6] + src.y * cXYZ_D65f[7] + src.z * cXYZ_D65f[8]; - } - - template struct XYZ2RGBBase - { - typedef int coeff_t; - - explicit XYZ2RGBBase(const coeff_t coeffs[9]) - { - cudaSafeCall( cudaMemcpyToSymbol(cXYZ_D65i, coeffs, 9 * sizeof(int)) ); - } - }; - template <> struct XYZ2RGBBase - { - typedef float coeff_t; - - explicit XYZ2RGBBase(const coeff_t coeffs[9]) - { - cudaSafeCall( cudaMemcpyToSymbol(cXYZ_D65f, coeffs, 9 * sizeof(float)) ); - } - }; - template struct XYZ2RGB : XYZ2RGBBase - { - typedef typename RGB2XYZBase::coeff_t coeff_t; - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - explicit XYZ2RGB(const coeff_t coeffs[9]) : XYZ2RGBBase(coeffs) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - XYZ2RGBConvert(src, &dst.x); - setAlpha(dst, ColorChannel::max()); - return dst; - } - }; - - template - void RGB2XYZ_caller(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream) - { - typedef typename RGB2XYZ::coeff_t coeff_t; - RGB2XYZ cvt((const coeff_t*)coeffs); - callConvert(src, dst, cvt, stream); - } - - void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = - { - {RGB2XYZ_caller, RGB2XYZ_caller}, - {RGB2XYZ_caller, RGB2XYZ_caller} - }; - - RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - - void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = - { - {RGB2XYZ_caller, RGB2XYZ_caller}, - {RGB2XYZ_caller, RGB2XYZ_caller} - }; - - RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - - void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = - { - {RGB2XYZ_caller, RGB2XYZ_caller}, - {RGB2XYZ_caller, RGB2XYZ_caller} - }; - - RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - - template - void XYZ2RGB_caller(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream) - { - typedef typename XYZ2RGB::coeff_t coeff_t; - XYZ2RGB cvt((const coeff_t*)coeffs); - callConvert(src, dst, cvt, stream); - } - - void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = - { - {XYZ2RGB_caller, XYZ2RGB_caller}, - {XYZ2RGB_caller, XYZ2RGB_caller} - }; - - XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - - void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = - { - {XYZ2RGB_caller, XYZ2RGB_caller}, - {XYZ2RGB_caller, XYZ2RGB_caller} - }; - - XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - - void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const void* coeffs, cudaStream_t stream) - { - typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, const void* coeffs, cudaStream_t stream); - static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = - { - {XYZ2RGB_caller, XYZ2RGB_caller}, - {XYZ2RGB_caller, XYZ2RGB_caller} - }; - - XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, coeffs, stream); - } - -////////////////////////////////////// RGB <-> HSV /////////////////////////////////////// - - __constant__ int cHsvDivTable [256] = {0, 1044480, 522240, 348160, 261120, 208896, 174080, 149211, 130560, 116053, 104448, 94953, 87040, 80345, 74606, 69632, 65280, 61440, 58027, 54973, 52224, 49737, 47476, 45412, 43520, 41779, 40172, 38684, 37303, 36017, 34816, 33693, 32640, 31651, 30720, 29842, 29013, 28229, 27486, 26782, 26112, 25475, 24869, 24290, 23738, 23211, 22706, 22223, 21760, 21316, 20890, 20480, 20086, 19707, 19342, 18991, 18651, 18324, 18008, 17703, 17408, 17123, 16846, 16579, 16320, 16069, 15825, 15589, 15360, 15137, 14921, 14711, 14507, 14308, 14115, 13926, 13743, 13565, 13391, 13221, 13056, 12895, 12738, 12584, 12434, 12288, 12145, 12006, 11869, 11736, 11605, 11478, 11353, 11231, 11111, 10995, 10880, 10768, 10658, 10550, 10445, 10341, 10240, 10141, 10043, 9947, 9854, 9761, 9671, 9582, 9495, 9410, 9326, 9243, 9162, 9082, 9004, 8927, 8852, 8777, 8704, 8632, 8561, 8492, 8423, 8356, 8290, 8224, 8160, 8097, 8034, 7973, 7913, 7853, 7795, 7737, 7680, 7624, 7569, 7514, 7461, 7408, 7355, 7304, 7253, 7203, 7154, 7105, 7057, 7010, 6963, 6917, 6872, 6827, 6782, 6739, 6695, 6653, 6611, 6569, 6528, 6487, 6447, 6408, 6369, 6330, 6292, 6254, 6217, 6180, 6144, 6108, 6073, 6037, 6003, 5968, 5935, 5901, 5868, 5835, 5803, 5771, 5739, 5708, 5677, 5646, 5615, 5585, 5556, 5526, 5497, 5468, 5440, 5412, 5384, 5356, 5329, 5302, 5275, 5249, 5222, 5196, 5171, 5145, 5120, 5095, 5070, 5046, 5022, 4998, 4974, 4950, 4927, 4904, 4881, 4858, 4836, 4813, 4791, 4769, 4748, 4726, 4705, 4684, 4663, 4642, 4622, 4601, 4581, 4561, 4541, 4522, 4502, 4483, 4464, 4445, 4426, 4407, 4389, 4370, 4352, 4334, 4316, 4298, 4281, 4263, 4246, 4229, 4212, 4195, 4178, 4161, 4145, 4128, 4112, 4096}; - __constant__ int cHsvDivTable180[256] = {0, 122880, 61440, 40960, 30720, 24576, 20480, 17554, 15360, 13653, 12288, 11171, 10240, 9452, 8777, 8192, 7680, 7228, 6827, 6467, 6144, 5851, 5585, 5343, 5120, 4915, 4726, 4551, 4389, 4237, 4096, 3964, 3840, 3724, 3614, 3511, 3413, 3321, 3234, 3151, 3072, 2997, 2926, 2858, 2793, 2731, 2671, 2614, 2560, 2508, 2458, 2409, 2363, 2318, 2276, 2234, 2194, 2156, 2119, 2083, 2048, 2014, 1982, 1950, 1920, 1890, 1862, 1834, 1807, 1781, 1755, 1731, 1707, 1683, 1661, 1638, 1617, 1596, 1575, 1555, 1536, 1517, 1499, 1480, 1463, 1446, 1429, 1412, 1396, 1381, 1365, 1350, 1336, 1321, 1307, 1293, 1280, 1267, 1254, 1241, 1229, 1217, 1205, 1193, 1182, 1170, 1159, 1148, 1138, 1127, 1117, 1107, 1097, 1087, 1078, 1069, 1059, 1050, 1041, 1033, 1024, 1016, 1007, 999, 991, 983, 975, 968, 960, 953, 945, 938, 931, 924, 917, 910, 904, 897, 890, 884, 878, 871, 865, 859, 853, 847, 842, 836, 830, 825, 819, 814, 808, 803, 798, 793, 788, 783, 778, 773, 768, 763, 759, 754, 749, 745, 740, 736, 731, 727, 723, 719, 714, 710, 706, 702, 698, 694, 690, 686, 683, 679, 675, 671, 668, 664, 661, 657, 654, 650, 647, 643, 640, 637, 633, 630, 627, 624, 621, 617, 614, 611, 608, 605, 602, 599, 597, 594, 591, 588, 585, 582, 580, 577, 574, 572, 569, 566, 564, 561, 559, 556, 554, 551, 549, 546, 544, 541, 539, 537, 534, 532, 530, 527, 525, 523, 521, 518, 516, 514, 512, 510, 508, 506, 504, 502, 500, 497, 495, 493, 492, 490, 488, 486, 484, 482}; - __constant__ int cHsvDivTable256[256] = {0, 174763, 87381, 58254, 43691, 34953, 29127, 24966, 21845, 19418, 17476, 15888, 14564, 13443, 12483, 11651, 10923, 10280, 9709, 9198, 8738, 8322, 7944, 7598, 7282, 6991, 6722, 6473, 6242, 6026, 5825, 5638, 5461, 5296, 5140, 4993, 4855, 4723, 4599, 4481, 4369, 4263, 4161, 4064, 3972, 3884, 3799, 3718, 3641, 3567, 3495, 3427, 3361, 3297, 3236, 3178, 3121, 3066, 3013, 2962, 2913, 2865, 2819, 2774, 2731, 2689, 2648, 2608, 2570, 2533, 2497, 2461, 2427, 2394, 2362, 2330, 2300, 2270, 2241, 2212, 2185, 2158, 2131, 2106, 2081, 2056, 2032, 2009, 1986, 1964, 1942, 1920, 1900, 1879, 1859, 1840, 1820, 1802, 1783, 1765, 1748, 1730, 1713, 1697, 1680, 1664, 1649, 1633, 1618, 1603, 1589, 1574, 1560, 1547, 1533, 1520, 1507, 1494, 1481, 1469, 1456, 1444, 1432, 1421, 1409, 1398, 1387, 1376, 1365, 1355, 1344, 1334, 1324, 1314, 1304, 1295, 1285, 1276, 1266, 1257, 1248, 1239, 1231, 1222, 1214, 1205, 1197, 1189, 1181, 1173, 1165, 1157, 1150, 1142, 1135, 1128, 1120, 1113, 1106, 1099, 1092, 1085, 1079, 1072, 1066, 1059, 1053, 1046, 1040, 1034, 1028, 1022, 1016, 1010, 1004, 999, 993, 987, 982, 976, 971, 966, 960, 955, 950, 945, 940, 935, 930, 925, 920, 915, 910, 906, 901, 896, 892, 887, 883, 878, 874, 869, 865, 861, 857, 853, 848, 844, 840, 836, 832, 828, 824, 820, 817, 813, 809, 805, 802, 798, 794, 791, 787, 784, 780, 777, 773, 770, 767, 763, 760, 757, 753, 750, 747, 744, 741, 737, 734, 731, 728, 725, 722, 719, 716, 713, 710, 708, 705, 702, 699, 696, 694, 691, 688, 685}; - - template __device__ void RGB2HSVConvert(const uchar* src, D& dst, int bidx, int hr) - { - const int hsv_shift = 12; - const int* hdiv_table = hr == 180 ? cHsvDivTable180 : cHsvDivTable256; - - int b = src[bidx], g = src[1], r = src[bidx^2]; - int h, s, v = b; - int vmin = b, diff; - int vr, vg; - - v = max(v, g); - v = max(v, r); - vmin = min(vmin, g); - vmin = min(vmin, r); - - diff = v - vmin; - vr = v == r ? -1 : 0; - vg = v == g ? -1 : 0; - - s = (diff * cHsvDivTable[v] + (1 << (hsv_shift-1))) >> hsv_shift; - h = (vr & (g - b)) + (~vr & ((vg & (b - r + 2 * diff)) + ((~vg) & (r - g + 4 * diff)))); - h = (h * hdiv_table[diff] + (1 << (hsv_shift-1))) >> hsv_shift; - h += h < 0 ? hr : 0; - - dst.x = saturate_cast(h); - dst.y = (uchar)s; - dst.z = (uchar)v; - } - template __device__ void RGB2HSVConvert(const float* src, D& dst, int bidx, int hr) - { - const float hscale = hr * (1.f / 360.f); - - float b = src[bidx], g = src[1], r = src[bidx^2]; - float h, s, v; - - float vmin, diff; - - v = vmin = r; - v = fmax(v, g); - v = fmax(v, b); - vmin = fmin(vmin, g); - vmin = fmin(vmin, b); - - diff = v - vmin; - s = diff / (float)(fabs(v) + numeric_limits_gpu::epsilon()); - diff = (float)(60. / (diff + numeric_limits_gpu::epsilon())); - - if (v == r) - h = (g - b) * diff; - else if (v == g) - h = (b - r) * diff + 120.f; - else - h = (r - g) * diff + 240.f; - - if (h < 0) h += 360.f; - - dst.x = h * hscale; - dst.y = s; - dst.z = v; - } - - template struct RGB2HSV - { - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - RGB2HSV(int bidx, int hr) : bidx(bidx), hr(hr) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - RGB2HSVConvert(&src.x, dst, bidx, hr); - return dst; - } - - private: - int bidx; - int hr; - }; - - __constant__ int cHsvSectorData[6][3] = - { - {1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0} - }; - - template __device__ void HSV2RGBConvert(const T& src, float* dst, int bidx, int hr) - { - const float hscale = 6.f / hr; - - float h = src.x, s = src.y, v = src.z; - float b, g, r; - - if( s == 0 ) - b = g = r = v; - else - { - float tab[4]; - int sector; - h *= hscale; - if( h < 0 ) - do h += 6; while( h < 0 ); - else if( h >= 6 ) - do h -= 6; while( h >= 6 ); - sector = __float2int_rd(h); - h -= sector; - - tab[0] = v; - tab[1] = v*(1.f - s); - tab[2] = v*(1.f - s*h); - tab[3] = v*(1.f - s*(1.f - h)); - - b = tab[cHsvSectorData[sector][0]]; - g = tab[cHsvSectorData[sector][1]]; - r = tab[cHsvSectorData[sector][2]]; - } - - dst[bidx] = b; - dst[1] = g; - dst[bidx^2] = r; - } - template __device__ void HSV2RGBConvert(const T& src, uchar* dst, int bidx, int hr) - { - float3 buf; - - buf.x = src.x; - buf.y = src.y * (1.f/255.f); - buf.z = src.z * (1.f/255.f); - - HSV2RGBConvert(buf, &buf.x, bidx, hr); - - dst[0] = saturate_cast(buf.x * 255.f); - dst[1] = saturate_cast(buf.y * 255.f); - dst[2] = saturate_cast(buf.z * 255.f); - } - - template struct HSV2RGB - { - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - HSV2RGB(int bidx, int hr) : bidx(bidx), hr(hr) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - HSV2RGBConvert(src, &dst.x, bidx, hr); - setAlpha(dst, ColorChannel::max()); - return dst; - } - - private: - int bidx; - int hr; - }; - - template - void RGB2HSV_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream) - { - RGB2HSV cvt(bidx, hrange); - callConvert(src, dst, cvt, stream); - } - - void RGB2HSV_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*RGB2HSV_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const RGB2HSV_caller_t RGB2HSV_callers[2][2] = - { - {RGB2HSV_caller, RGB2HSV_caller}, - {RGB2HSV_caller, RGB2HSV_caller} - }; - - RGB2HSV_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - void RGB2HSV_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*RGB2HSV_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const RGB2HSV_caller_t RGB2HSV_callers[2][2] = - { - {RGB2HSV_caller, RGB2HSV_caller}, - {RGB2HSV_caller, RGB2HSV_caller} - }; - - RGB2HSV_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - template - void HSV2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream) - { - HSV2RGB cvt(bidx, hrange); - callConvert(src, dst, cvt, stream); - } - - void HSV2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*HSV2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const HSV2RGB_caller_t HSV2RGB_callers[2][2] = - { - {HSV2RGB_caller, HSV2RGB_caller}, - {HSV2RGB_caller, HSV2RGB_caller} - }; - - HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - void HSV2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*HSV2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const HSV2RGB_caller_t HSV2RGB_callers[2][2] = - { - {HSV2RGB_caller, HSV2RGB_caller}, - {HSV2RGB_caller, HSV2RGB_caller} - }; - - HSV2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - -/////////////////////////////////////// RGB <-> HLS //////////////////////////////////////// - - template __device__ void RGB2HLSConvert(const float* src, D& dst, int bidx, int hr) - { - const float hscale = hr * (1.f / 360.f); - - float b = src[bidx], g = src[1], r = src[bidx^2]; - float h = 0.f, s = 0.f, l; - float vmin, vmax, diff; - - vmax = vmin = r; - vmax = fmax(vmax, g); - vmax = fmax(vmax, b); - vmin = fmin(vmin, g); - vmin = fmin(vmin, b); - - diff = vmax - vmin; - l = (vmax + vmin) * 0.5f; - - if (diff > numeric_limits_gpu::epsilon()) - { - s = l < 0.5f ? diff / (vmax + vmin) : diff / (2.0f - vmax - vmin); - diff = 60.f / diff; - - if (vmax == r) - h = (g - b)*diff; - else if (vmax == g) - h = (b - r)*diff + 120.f; - else - h = (r - g)*diff + 240.f; - - if (h < 0.f) h += 360.f; - } - - dst.x = h * hscale; - dst.y = l; - dst.z = s; - } - template __device__ void RGB2HLSConvert(const uchar* src, D& dst, int bidx, int hr) - { - float3 buf; - - buf.x = src[0]*(1.f/255.f); - buf.y = src[1]*(1.f/255.f); - buf.z = src[2]*(1.f/255.f); - - RGB2HLSConvert(&buf.x, buf, bidx, hr); - - dst.x = saturate_cast(buf.x); - dst.y = saturate_cast(buf.y*255.f); - dst.z = saturate_cast(buf.z*255.f); - } - - template struct RGB2HLS - { - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - RGB2HLS(int bidx, int hr) : bidx(bidx), hr(hr) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - RGB2HLSConvert(&src.x, dst, bidx, hr); - return dst; - } - - private: - int bidx; - int hr; - }; - - __constant__ int cHlsSectorData[6][3] = - { - {1,3,0}, {1,0,2}, {3,0,1}, {0,2,1}, {0,1,3}, {2,1,0} - }; - - template __device__ void HLS2RGBConvert(const T& src, float* dst, int bidx, int hr) - { - const float hscale = 6.0f / hr; - - float h = src.x, l = src.y, s = src.z; - float b, g, r; - - if (s == 0) - b = g = r = l; - else - { - float tab[4]; - int sector; - - float p2 = l <= 0.5f ? l * (1 + s) : l + s - l * s; - float p1 = 2 * l - p2; - - h *= hscale; - - if( h < 0 ) - do h += 6; while( h < 0 ); - else if( h >= 6 ) - do h -= 6; while( h >= 6 ); - - sector = __float2int_rd(h); - h -= sector; - - tab[0] = p2; - tab[1] = p1; - tab[2] = p1 + (p2 - p1) * (1 - h); - tab[3] = p1 + (p2 - p1) * h; - - b = tab[cHlsSectorData[sector][0]]; - g = tab[cHlsSectorData[sector][1]]; - r = tab[cHlsSectorData[sector][2]]; - } - - dst[bidx] = b; - dst[1] = g; - dst[bidx^2] = r; - } - template __device__ void HLS2RGBConvert(const T& src, uchar* dst, int bidx, int hr) - { - float3 buf; - - buf.x = src.x; - buf.y = src.y*(1.f/255.f); - buf.z = src.z*(1.f/255.f); - - HLS2RGBConvert(buf, &buf.x, bidx, hr); - - dst[0] = saturate_cast(buf.x*255.f); - dst[1] = saturate_cast(buf.y*255.f); - dst[2] = saturate_cast(buf.z*255.f); - } - - template struct HLS2RGB - { - typedef typename TypeVec::vec_t src_t; - typedef typename TypeVec::vec_t dst_t; - - HLS2RGB(int bidx, int hr) : bidx(bidx), hr(hr) {} - - __device__ __forceinline__ dst_t operator()(const src_t& src) const - { - dst_t dst; - HLS2RGBConvert(src, &dst.x, bidx, hr); - setAlpha(dst, ColorChannel::max()); - return dst; - } - - private: - int bidx; - int hr; - }; - - template - void RGB2HLS_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream) - { - RGB2HLS cvt(bidx, hrange); - callConvert(src, dst, cvt, stream); - } - - void RGB2HLS_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*RGB2HLS_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const RGB2HLS_caller_t RGB2HLS_callers[2][2] = - { - {RGB2HLS_caller, RGB2HLS_caller}, - {RGB2HLS_caller, RGB2HLS_caller} - }; - - RGB2HLS_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - void RGB2HLS_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*RGB2HLS_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const RGB2HLS_caller_t RGB2HLS_callers[2][2] = - { - {RGB2HLS_caller, RGB2HLS_caller}, - {RGB2HLS_caller, RGB2HLS_caller} - }; - - RGB2HLS_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - - template - void HLS2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream) - { - HLS2RGB cvt(bidx, hrange); - callConvert(src, dst, cvt, stream); - } - - void HLS2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*HLS2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const HLS2RGB_caller_t HLS2RGB_callers[2][2] = - { - {HLS2RGB_caller, HLS2RGB_caller}, - {HLS2RGB_caller, HLS2RGB_caller} - }; - - HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } - - void HLS2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, int hrange, cudaStream_t stream) - { - typedef void (*HLS2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, int hrange, cudaStream_t stream); - static const HLS2RGB_caller_t HLS2RGB_callers[2][2] = - { - {HLS2RGB_caller, HLS2RGB_caller}, - {HLS2RGB_caller, HLS2RGB_caller} - }; - - HLS2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, hrange, stream); - } + #define OPENCV_GPU_IMPLEMENT_CVTCOLOR(name, traits) \ + void name(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) \ + { \ + traits::functor_type functor = traits::create_functor(); \ + typedef typename traits::functor_type::argument_type src_t; \ + typedef typename traits::functor_type::result_type dst_t; \ + transform((DevMem2D_)src, (DevMem2D_)dst, functor, stream); \ + } + + #define OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(name) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name, name ## _traits) + + #define OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(name) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _8u, name ## _traits) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _16u, name ## _traits) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _32f, name ## _traits) + + #define OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(name) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _8u, name ## _traits) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _32f, name ## _traits) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _full_8u, name ## _full_traits) \ + OPENCV_GPU_IMPLEMENT_CVTCOLOR(name ## _full_32f, name ## _full_traits) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_rgba) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr_to_bgr555) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr_to_bgr565) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(rgb_to_bgr555) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(rgb_to_bgr565) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgra_to_bgr555) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgra_to_bgr565) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(rgba_to_bgr555) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(rgba_to_bgr565) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr555_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr565_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr555_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr565_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr555_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr565_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr555_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr565_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(gray_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(gray_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(gray_to_bgr555) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(gray_to_bgr565) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr555_to_gray) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE(bgr565_to_gray) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_gray) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_gray) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_gray) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_gray) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_yuv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_yuv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_yuv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_yuv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_yuv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_yuv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_yuv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_yuv4) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv4_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv4_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv4_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(yuv4_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_YCrCb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_YCrCb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_YCrCb4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_YCrCb4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_YCrCb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_YCrCb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_YCrCb4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_YCrCb4) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb4_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb4_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb4_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(YCrCb4_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_xyz) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_xyz) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgb_to_xyz4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(rgba_to_xyz4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_xyz) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_xyz) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgr_to_xyz4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(bgra_to_xyz4) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz4_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz4_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz4_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL(xyz4_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgb_to_hsv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgba_to_hsv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgb_to_hsv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgba_to_hsv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgr_to_hsv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgra_to_hsv) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgr_to_hsv4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgra_to_hsv4) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv4_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv4_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv4_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hsv4_to_bgra) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgb_to_hls) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgba_to_hls) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgb_to_hls4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(rgba_to_hls4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgr_to_hls) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgra_to_hls) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgr_to_hls4) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(bgra_to_hls4) + + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls4_to_rgb) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls4_to_rgba) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls_to_bgra) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls4_to_bgr) + OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F(hls4_to_bgra) + + #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR + #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_ONE + #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_ALL + #undef OPENCV_GPU_IMPLEMENT_CVTCOLOR_8U32F }}} diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu index 84345547ed..69020dd22f 100644 --- a/modules/gpu/src/cuda/element_operations.cu +++ b/modules/gpu/src/cuda/element_operations.cu @@ -40,9 +40,10 @@ // //M*/ -#include "opencv2/gpu/device/vecmath.hpp" +#include "opencv2/gpu/device/functional.hpp" +#include "opencv2/gpu/device/vec_math.hpp" #include "opencv2/gpu/device/transform.hpp" -#include "opencv2/gpu/device/limits_gpu.hpp" +#include "opencv2/gpu/device/limits.hpp" #include "opencv2/gpu/device/saturate_cast.hpp" #include "internal_shared.hpp" @@ -354,114 +355,11 @@ namespace cv { namespace gpu { namespace mathfunc ////////////////////////////////////////////////////////////////////////// // min/max - - struct MinOp - { - template - __device__ __forceinline__ T operator()(T a, T b) - { - return min(a, b); - } - __device__ __forceinline__ float operator()(float a, float b) - { - return fmin(a, b); - } - __device__ __forceinline__ double operator()(double a, double b) - { - return fmin(a, b); - } - }; - - struct MaxOp - { - template - __device__ __forceinline__ T operator()(T a, T b) - { - return max(a, b); - } - __device__ __forceinline__ float operator()(float a, float b) - { - return fmax(a, b); - } - __device__ __forceinline__ double operator()(double a, double b) - { - return fmax(a, b); - } - }; - - template struct ScalarMinOp - { - T s; - - explicit ScalarMinOp(T s_) : s(s_) {} - - __device__ __forceinline__ T operator()(T a) - { - return min(a, s); - } - }; - template <> struct ScalarMinOp - { - float s; - - explicit ScalarMinOp(float s_) : s(s_) {} - - __device__ __forceinline__ float operator()(float a) - { - return fmin(a, s); - } - }; - template <> struct ScalarMinOp - { - double s; - - explicit ScalarMinOp(double s_) : s(s_) {} - - __device__ __forceinline__ double operator()(double a) - { - return fmin(a, s); - } - }; - - template struct ScalarMaxOp - { - T s; - - explicit ScalarMaxOp(T s_) : s(s_) {} - - __device__ __forceinline__ T operator()(T a) - { - return max(a, s); - } - }; - template <> struct ScalarMaxOp - { - float s; - - explicit ScalarMaxOp(float s_) : s(s_) {} - - __device__ __forceinline__ float operator()(float a) - { - return fmax(a, s); - } - }; - template <> struct ScalarMaxOp - { - double s; - - explicit ScalarMaxOp(double s_) : s(s_) {} - - __device__ __forceinline__ double operator()(double a) - { - return fmax(a, s); - } - }; template void min_gpu(const DevMem2D_& src1, const DevMem2D_& src2, const DevMem2D_& dst, cudaStream_t stream) { - MinOp op; - transform(src1, src2, dst, op, stream); + transform(src1, src2, dst, minimum(), stream); } template void min_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream); @@ -475,8 +373,7 @@ namespace cv { namespace gpu { namespace mathfunc template void max_gpu(const DevMem2D_& src1, const DevMem2D_& src2, const DevMem2D_& dst, cudaStream_t stream) { - MaxOp op; - transform(src1, src2, dst, op, stream); + transform(src1, src2, dst, maximum(), stream); } template void max_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream); @@ -490,8 +387,7 @@ namespace cv { namespace gpu { namespace mathfunc template void min_gpu(const DevMem2D_& src1, T src2, const DevMem2D_& dst, cudaStream_t stream) { - ScalarMinOp op(src2); - transform(src1, dst, op, stream); + transform(src1, dst, device::bind2nd(minimum(), src2), stream); } template void min_gpu(const DevMem2D& src1, uchar src2, const DevMem2D& dst, cudaStream_t stream); @@ -501,12 +397,11 @@ namespace cv { namespace gpu { namespace mathfunc template void min_gpu(const DevMem2D_& src1, int src2, const DevMem2D_& dst, cudaStream_t stream); template void min_gpu(const DevMem2D_& src1, float src2, const DevMem2D_& dst, cudaStream_t stream); template void min_gpu(const DevMem2D_& src1, double src2, const DevMem2D_& dst, cudaStream_t stream); - + template void max_gpu(const DevMem2D_& src1, T src2, const DevMem2D_& dst, cudaStream_t stream) { - ScalarMaxOp op(src2); - transform(src1, dst, op, stream); + transform(src1, dst, device::bind2nd(maximum(), src2), stream); } template void max_gpu(const DevMem2D& src1, uchar src2, const DevMem2D& dst, cudaStream_t stream); @@ -519,100 +414,7 @@ namespace cv { namespace gpu { namespace mathfunc ////////////////////////////////////////////////////////////////////////// - // threshold - - template struct ThreshBinary - { - ThreshBinary(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} - - __device__ __forceinline__ T operator()(const T& src) const - { - return src > thresh ? maxVal : 0; - } - - private: - T thresh; - T maxVal; - }; - - template struct ThreshBinaryInv - { - ThreshBinaryInv(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} - - __device__ __forceinline__ T operator()(const T& src) const - { - return src > thresh ? 0 : maxVal; - } - - private: - T thresh; - T maxVal; - }; - - template struct ThreshTrunc - { - ThreshTrunc(T thresh_, T) : thresh(thresh_) {} - - __device__ __forceinline__ T operator()(const T& src) const - { - return min(src, thresh); - } - - private: - T thresh; - }; - template <> struct ThreshTrunc - { - ThreshTrunc(float thresh_, float) : thresh(thresh_) {} - - __device__ __forceinline__ float operator()(const float& src) const - { - return fmin(src, thresh); - } - - private: - float thresh; - }; - template <> struct ThreshTrunc - { - ThreshTrunc(double thresh_, double) : thresh(thresh_) {} - - __device__ __forceinline__ double operator()(const double& src) const - { - return fmin(src, thresh); - } - - private: - double thresh; - }; - - template struct ThreshToZero - { - public: - ThreshToZero(T thresh_, T) : thresh(thresh_) {} - - __device__ __forceinline__ T operator()(const T& src) const - { - return src > thresh ? src : 0; - } - - private: - T thresh; - }; - - template struct ThreshToZeroInv - { - public: - ThreshToZeroInv(T thresh_, T) : thresh(thresh_) {} - - __device__ __forceinline__ T operator()(const T& src) const - { - return src > thresh ? 0 : src; - } - - private: - T thresh; - }; + // threshold template