Skip to content

Commit

Permalink
Extend cudaimgproc::demosaicing for f32
Browse files Browse the repository at this point in the history
  • Loading branch information
@apbr
apbr committed Nov 28, 2024
1 parent 5409e01 commit ec24b9c
Show file tree
Hide file tree
Showing 2 changed files with 61 additions and 16 deletions.
33 changes: 24 additions & 9 deletions modules/cudaimgproc/src/color.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -72,6 +72,8 @@ namespace cv { namespace cuda {

template <int cn>
void MHCdemosaic(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);

void MHCdemosaic_float3(PtrStepSzf src, int2 sourceOffset, PtrStepSzf dst, int2 firstRed, cudaStream_t stream);
}
}}

Expand Down Expand Up @@ -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 );
CV_Assert( depth == CV_8U || (depth == CV_32F && dcn == 3) );
CV_Assert( src.channels() == 1 );
CV_Assert( dcn == 3 || dcn == 4 );

Expand All @@ -2145,18 +2147,31 @@ void cv::cuda::demosaicing(InputArray _src, OutputArray _dst, int code, int dcn,

dst.setTo(Scalar::all(0), stream);

Size wholeSize;
Point ofs;
src.locateROI(wholeSize, ofs);
PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);

const int2 firstRed = make_int2(code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGB2BGR_MHT ? 0 : 1,
code == COLOR_BayerRG2BGR_MHT || code == COLOR_BayerGR2BGR_MHT ? 0 : 1);

if (dcn == 3)
cv::cuda::device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
if (depth == CV_8U)
{
Size wholeSize;
Point ofs;
src.locateROI(wholeSize, ofs);
PtrStepSzb srcWhole(wholeSize.height, wholeSize.width, src.datastart, src.step);

if (dcn == 3)
cv::cuda::device::MHCdemosaic<3>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
else
cv::cuda::device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
}
// depth == CV_32F && dcn == 3
else
cv::cuda::device::MHCdemosaic<4>(srcWhole, make_int2(ofs.x, ofs.y), dst, firstRed, StreamAccessor::getStream(stream));
{
Size wholeSize;
Point ofs;
src.locateROI(wholeSize, ofs);
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));
}

break;
}
Expand Down
44 changes: 37 additions & 7 deletions modules/cudaimgproc/src/cuda/debayer.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -382,6 +382,14 @@ 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 D> __device__ __forceinline__ D toDstColor(const float3& pix){
return toDst<D>(make_uchar3(saturate_cast<uchar>(pix.x), saturate_cast<uchar>(pix.y), saturate_cast<uchar>(pix.z)));
}
template <> __device__ __forceinline__ float3 toDstColor<float3>(const float3& pix)
{
return pix;
}

//////////////////////////////////////////////////////////////
// Bayer Demosaicing (Malvar, He, and Cutler)
//
Expand Down Expand Up @@ -507,16 +515,15 @@ namespace cv { namespace cuda { namespace device
alternate.y = (y + firstRed.y) % 2;

// in BGR sequence;
uchar3 pixelColor =
float3 pixelColor =
(alternate.y == 0) ?
((alternate.x == 0) ?
make_uchar3(saturate_cast<uchar>(PATTERN.y), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(C)) :
make_uchar3(saturate_cast<uchar>(PATTERN.w), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.z))) :
make_float3(PATTERN.y, PATTERN.x, C) :
make_float3(PATTERN.w, C, PATTERN.z)) :
((alternate.x == 0) ?
make_uchar3(saturate_cast<uchar>(PATTERN.z), saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.w)) :
make_uchar3(saturate_cast<uchar>(C), saturate_cast<uchar>(PATTERN.x), saturate_cast<uchar>(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>
Expand Down Expand Up @@ -545,6 +552,29 @@ namespace cv { namespace cuda { namespace device
template void MHCdemosaic<1>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
template void MHCdemosaic<3>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb dst, int2 firstRed, cudaStream_t stream);
template void MHCdemosaic<4>(PtrStepSzb src, int2 sourceOffset, PtrStepSzb 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 */

0 comments on commit ec24b9c

Please sign in to comment.