From d498de94c9d1dfbf855a93c9e3f394c7e891fa46 Mon Sep 17 00:00:00 2001 From: Jonathan Schnitzler Date: Thu, 28 Nov 2024 10:35:06 +0100 Subject: [PATCH] Extend cudaimgproc::demosaicing for f32 --- modules/cudaimgproc/src/color.cpp | 7 +++- modules/cudaimgproc/src/cuda/debayer.cu | 48 +++++++++++++++++++------ 2 files changed, 44 insertions(+), 11 deletions(-) diff --git a/modules/cudaimgproc/src/color.cpp b/modules/cudaimgproc/src/color.cpp index 94ffe90fa2f..b01b8d66fef 100644 --- a/modules/cudaimgproc/src/color.cpp +++ b/modules/cudaimgproc/src/color.cpp @@ -72,6 +72,8 @@ namespace cv { namespace cuda { template void MHCdemosaic(PtrStepSz src, int2 sourceOffset, PtrStepSz dst, int2 firstRed, cudaStream_t stream); + + void MHCdemosaic_float3(PtrStepSzf src, int2 sourceOffset, PtrStepSzf dst, int2 firstRed, cudaStream_t stream); } }} @@ -2136,7 +2138,7 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn, GpuMat src = _src.getGpuMat(); const int depth = _src.depth(); - CV_Assert( depth == CV_8U || depth == CV_16U); + CV_Assert( depth == CV_8U || depth == CV_16U || (depth == CV_32F && dcn == 3) ); CV_Assert( src.channels() == 1 ); CV_Assert( dcn == 3 || dcn == 4 ); @@ -2156,6 +2158,9 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn, if (depth == CV_8U) { PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step); cv::cuda::device::MHCdemosaic<3, uchar>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); + } else if (depth === CV_32F) { + PtrStepSzf srcWhole(wholeSize.height, wholeSize.width, (float*)src.datastart, src.step); + cv::cuda::device::MHCdemosaic_float3(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); } else { PtrStepSz srcWhole(wholeSize.height, wholeSize.width, src.ptr(), src.step); cv::cuda::device::MHCdemosaic<3, ushort>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream)); diff --git a/modules/cudaimgproc/src/cuda/debayer.cu b/modules/cudaimgproc/src/cuda/debayer.cu index dfd3b9aa11d..cf058c0f5b7 100644 --- a/modules/cudaimgproc/src/cuda/debayer.cu +++ b/modules/cudaimgproc/src/cuda/debayer.cu @@ -382,6 +382,15 @@ namespace cv { namespace cuda { namespace device template void Bayer2BGR_16u_gpu<3>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); template void Bayer2BGR_16u_gpu<4>(PtrStepSzb src, PtrStepSzb dst, bool blue_last, bool start_with_green, cudaStream_t stream); + template __device__ __forceinline__ DstType toDstColor(const float3& pix){ + typedef typename VecTraits::elem_type SrcElemType; + return toDst(make_3(saturate_cast(pix.x), saturate_cast(pix.y), saturate_cast(pix.z))); + } + template <> __device__ __forceinline__ float3 toDstColor(const float3& pix) + { + return pix; + } + ////////////////////////////////////////////////////////////// // Bayer Demosaicing (Malvar, He, and Cutler) // @@ -517,19 +526,15 @@ namespace cv { namespace cuda { namespace device alternate.x = (x + firstRed.x) % 2; alternate.y = (y + firstRed.y) % 2; - typedef typename VecTraits::elem_type SrcElemType; - typedef typename TypeVec::vec_type SrcType; - - SrcType pixelColor = + float3 pixelColor = (alternate.y == 0) ? ((alternate.x == 0) ? - make_3(saturate_cast(PATTERN.y), saturate_cast(PATTERN.x), saturate_cast(C)) : - make_3(saturate_cast(PATTERN.w), saturate_cast(C), saturate_cast(PATTERN.z))) : + make_float3(PATTERN.y, PATTERN.x, C) : + make_float3(PATTERN.w, C, PATTERN.z)) : ((alternate.x == 0) ? - make_3(saturate_cast(PATTERN.z), saturate_cast(C), saturate_cast(PATTERN.w)) : - make_3(saturate_cast(C), saturate_cast(PATTERN.x), saturate_cast(PATTERN.y))); - - dst(y, x) = toDst(pixelColor); + make_float3(PATTERN.z, C, PATTERN.w) : + make_float3(C, PATTERN.x, PATTERN.y)); + dst(y, x) = toDstColor(pixelColor); } template @@ -561,6 +566,29 @@ namespace cv { namespace cuda { namespace device template void MHCdemosaic<1, ushort>(PtrStepSz src, int2 sourceOffset, PtrStepSz dst, int2 firstRed, cudaStream_t stream); template void MHCdemosaic<3, ushort>(PtrStepSz src, int2 sourceOffset, PtrStepSz dst, int2 firstRed, cudaStream_t stream); template void MHCdemosaic<4, ushort>(PtrStepSz src, int2 sourceOffset, PtrStepSz dst, int2 firstRed, cudaStream_t stream); + + // Implement MHCdemosaic for float and with a result of 3 channels + void MHCdemosaic_float3(PtrStepSzf src, int2 sourceOffset, PtrStepSzf dst, int2 firstRed, cudaStream_t stream) + { + typedef typename TypeVec::vec_type dst_t; + + const dim3 block(32, 8); + const dim3 grid(divUp(src.cols, block.x), divUp(src.rows, block.y)); + + if (sourceOffset.x || sourceOffset.y) { + cv::cudev::TextureOff texSrc(src, sourceOffset.y, sourceOffset.x); + MHCdemosaic><<>>((PtrStepSz)dst, texSrc, firstRed); + } + else { + cv::cudev::Texture texSrc(src); + MHCdemosaic><<>>((PtrStepSz)dst, texSrc, firstRed); + } + + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } }}} #endif /* CUDA_DISABLER */