Extend cudaimgproc::demosaicing for f32

Author: apbr

modules/cudaimgproc/src/color.cpp

@@ -72,6 +72,8 @@ namespace cv { namespace cuda {
 
         template <int cn, typename Depth>
         void MHCdemosaic(PtrStepSz<Depth> src, int2 sourceOffset, PtrStepSz<Depth> 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<ushort> srcWhole(wholeSize.height, wholeSize.width, src.ptr<ushort>(), src.step);
                 cv::cuda::device::MHCdemosaic<3, ushort>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));

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 <typename DstType> __device__ __forceinline__ DstType toDstColor(const float3& pix){
+        typedef typename VecTraits<DstType>::elem_type SrcElemType;
+        return toDst<D>(make_3<SrcElemType>(saturate_cast<SrcElemType>(pix.x), saturate_cast<SrcElemType>(pix.y), saturate_cast<SrcElemType>(pix.z)));
+    }
+    template <> __device__ __forceinline__ float3 toDstColor<float3>(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<DstType>::elem_type SrcElemType;
-        typedef typename TypeVec<SrcElemType, 3>::vec_type SrcType;
-
-        SrcType pixelColor =
+        float3 pixelColor =
             (alternate.y == 0) ?
                 ((alternate.x == 0) ?
-                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.y), saturate_cast<SrcElemType>(PATTERN.x), saturate_cast<SrcElemType>(C)) :
-                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.w), saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.z))) :
+                    make_float3(PATTERN.y, PATTERN.x, C) :
+                    make_float3(PATTERN.w, C, PATTERN.z)) :
                 ((alternate.x == 0) ?
-                    make_3<SrcElemType>(saturate_cast<SrcElemType>(PATTERN.z), saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.w)) :
-                    make_3<SrcElemType>(saturate_cast<SrcElemType>(C), saturate_cast<SrcElemType>(PATTERN.x), saturate_cast<SrcElemType>(PATTERN.y)));
-
-        dst(y, x) = toDst<DstType>(pixelColor);
+                    make_float3(PATTERN.z, C, PATTERN.w) :
+                    make_float3(C, PATTERN.x, PATTERN.y));
+        dst(y, x) = toDstColor<DstType>(pixelColor);
     }
 
     template <int cn, typename Depth>
@@ -561,6 +566,29 @@ namespace cv { namespace cuda { namespace device
     template void MHCdemosaic<1, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> dst, int2 firstRed, cudaStream_t stream);
     template void MHCdemosaic<3, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> dst, int2 firstRed, cudaStream_t stream);
     template void MHCdemosaic<4, ushort>(PtrStepSz<ushort> src, int2 sourceOffset, PtrStepSz<ushort> 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<float, 3>::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<float> texSrc(src, sourceOffset.y, sourceOffset.x);
+            MHCdemosaic<dst_t, cv::cudev::TextureOffPtr<float>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
+        }
+        else {
+            cv::cudev::Texture<float> texSrc(src);
+            MHCdemosaic<dst_t, cv::cudev::TexturePtr<float>><<<grid, block, 0, stream>>>((PtrStepSz<dst_t>)dst, texSrc, firstRed);
+        }
+
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
 }}}
 
 #endif /* CUDA_DISABLER */