From adf7cc20b7f301c4c0e79b66927746551ea1ab2d Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Mon, 11 Oct 2010 14:25:30 +0000 Subject: [PATCH] changed mean shift procedure a little --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 + modules/gpu/src/cuda/imgproc.cu | 139 ++++++++++++++++++++------------ modules/gpu/src/imgproc_gpu.cpp | 29 +++++++ 3 files changed, 121 insertions(+), 51 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 7d0074d..8bb7c2a 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -469,6 +469,10 @@ namespace cv CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); + //! Does mean shift procedure on GPU. + CV_EXPORTS void meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, + TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); + //! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. //! Supported types of input disparity: CV_8U, CV_16S. //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255). diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 1dd4616..938e2d1 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -163,70 +163,93 @@ namespace imgproc { texture tex_meanshift; - extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, - int sp, int sr, int maxIter, float eps ) + __device__ short2 do_mean_shift(int x0, int y0, unsigned char* out, + int out_step, int cols, int rows, + int sp, int sr, int maxIter, float eps) { - int x0 = blockIdx.x * blockDim.x + threadIdx.x; - int y0 = blockIdx.y * blockDim.y + threadIdx.y; + int isr2 = sr*sr; + uchar4 c = tex2D(tex_meanshift, x0, y0 ); - if( x0 < cols && y0 < rows ) + // iterate meanshift procedure + for( int iter = 0; iter < maxIter; iter++ ) { - int isr2 = sr*sr; - uchar4 c = tex2D(tex_meanshift, x0, y0 ); - // iterate meanshift procedure - for( int iter = 0; iter < maxIter; iter++ ) + int count = 0; + int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0; + float icount; + + //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp) + int minx = x0-sp; + int miny = y0-sp; + int maxx = x0+sp; + int maxy = y0+sp; + + for( int y = miny; y <= maxy; y++) { - int count = 0; - int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0; - float icount; - - //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp) - int minx = x0-sp; - int miny = y0-sp; - int maxx = x0+sp; - int maxy = y0+sp; - - for( int y = miny; y <= maxy; y++) - { - int rowCount = 0; - for( int x = minx; x <= maxx; x++ ) - { - uchar4 t = tex2D( tex_meanshift, x, y ); - - int norm2 = (t.x - c.x) * (t.x - c.x) + (t.y - c.y) * (t.y - c.y) + (t.z - c.z) * (t.z - c.z); - if( norm2 <= isr2 ) - { - s0 += t.x; s1 += t.y; s2 += t.z; - sx += x; rowCount++; - } + int rowCount = 0; + for( int x = minx; x <= maxx; x++ ) + { + uchar4 t = tex2D( tex_meanshift, x, y ); + + int norm2 = (t.x - c.x) * (t.x - c.x) + (t.y - c.y) * (t.y - c.y) + (t.z - c.z) * (t.z - c.z); + if( norm2 <= isr2 ) + { + s0 += t.x; s1 += t.y; s2 += t.z; + sx += x; rowCount++; } - count += rowCount; - sy += y*rowCount; } + count += rowCount; + sy += y*rowCount; + } - if( count == 0 ) - break; + if( count == 0 ) + break; - icount = 1.f/count; - int x1 = __float2int_rz(sx*icount); - int y1 = __float2int_rz(sy*icount); - s0 = __float2int_rz(s0*icount); - s1 = __float2int_rz(s1*icount); - s2 = __float2int_rz(s2*icount); + icount = 1.f/count; + int x1 = __float2int_rz(sx*icount); + int y1 = __float2int_rz(sy*icount); + s0 = __float2int_rz(s0*icount); + s1 = __float2int_rz(s1*icount); + s2 = __float2int_rz(s2*icount); - int norm2 = (s0 - c.x) * (s0 - c.x) + (s1 - c.y) * (s1 - c.y) + (s2 - c.z) * (s2 - c.z); + int norm2 = (s0 - c.x) * (s0 - c.x) + (s1 - c.y) * (s1 - c.y) + (s2 - c.z) * (s2 - c.z); - bool stopFlag = (x0 == x1 && y0 == y1) || (abs(x1-x0) + abs(y1-y0) + norm2 <= eps); + bool stopFlag = (x0 == x1 && y0 == y1) || (abs(x1-x0) + abs(y1-y0) + norm2 <= eps); - x0 = x1; y0 = y1; - c.x = s0; c.y = s1; c.z = s2; + x0 = x1; y0 = y1; + c.x = s0; c.y = s1; c.z = s2; - if( stopFlag ) - break; - } + if( stopFlag ) + break; + } + + int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar); + *(uchar4*)(out + base) = c; + + return make_short2((short)x0, (short)y0); + } + + extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, + int sp, int sr, int maxIter, float eps ) + { + int x0 = blockIdx.x * blockDim.x + threadIdx.x; + int y0 = blockIdx.y * blockDim.y + threadIdx.y; + + if( x0 < cols && y0 < rows ) + do_mean_shift(x0, y0, out, out_step, cols, rows, sp, sr, maxIter, eps); + } + + extern "C" __global__ void meanshiftproc_kernel( unsigned char* outr, int outrstep, + unsigned char* outsp, int outspstep, + int cols, int rows, + int sp, int sr, int maxIter, float eps ) + { + int x0 = blockIdx.x * blockDim.x + threadIdx.x; + int y0 = blockIdx.y * blockDim.y + threadIdx.y; - int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar); - *(uchar4*)(out + base) = c; + if( x0 < cols && y0 < rows ) + { + int basesp = (blockIdx.y * blockDim.y + threadIdx.y) * outspstep + (blockIdx.x * blockDim.x + threadIdx.x) * 2 * sizeof(short); + *(short2*)(outsp + basesp) = do_mean_shift(x0, y0, outr, outrstep, cols, rows, sp, sr, maxIter, eps); } } } @@ -247,6 +270,20 @@ namespace cv { namespace gpu { namespace improc cudaSafeCall( cudaThreadSynchronize() ); cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) ); } + extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps) + { + dim3 grid(1, 1, 1); + dim3 threads(32, 16, 1); + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); + + cudaChannelFormatDesc desc = cudaCreateChannelDesc(); + cudaSafeCall( cudaBindTexture2D( 0, imgproc::tex_meanshift, src.ptr, desc, src.cols, src.rows, src.step ) ); + + imgproc::meanshiftproc_kernel<<< grid, threads >>>( dstr.ptr, dstr.step, dstsp.ptr, dstsp.step, dstr.cols, dstr.rows, sp, sr, maxIter, eps ); + cudaSafeCall( cudaThreadSynchronize() ); + cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) ); + } }}} /////////////////////////////////// drawColorDisp /////////////////////////////////////////////// diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 0badf8b..33bf5b5 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -49,6 +49,7 @@ using namespace cv::gpu; void cv::gpu::remap(const GpuMat&, GpuMat&, const GpuMat&, const GpuMat&){ throw_nogpu(); } void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); } +void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); } void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); } void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); } void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); } @@ -74,6 +75,7 @@ namespace cv { namespace gpu void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst); extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps); + extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps); void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream); void drawColorDisp_gpu(const DevMem2D_& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream); @@ -164,6 +166,33 @@ void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, } //////////////////////////////////////////////////////////////////////// +// meanShiftProc_GPU + +void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, TermCriteria criteria) +{ + if( src.empty() ) + CV_Error( CV_StsBadArg, "The input image is empty" ); + + if( src.depth() != CV_8U || src.channels() != 4 ) + CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" ); + + dstr.create( src.size(), CV_8UC4 ); + dstsp.create( src.size(), CV_16SC2 ); + + if( !(criteria.type & TermCriteria::MAX_ITER) ) + criteria.maxCount = 5; + + int maxIter = std::min(std::max(criteria.maxCount, 1), 100); + + float eps; + if( !(criteria.type & TermCriteria::EPS) ) + eps = 1.f; + eps = (float)std::max(criteria.epsilon, 0.0); + + improc::meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps); +} + +//////////////////////////////////////////////////////////////////////// // drawColorDisp namespace -- 2.7.4