From e5a3ab3cb97d8a14619fd96cd23f7aa626eb2060 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 11 Jul 2014 15:01:46 +0400 Subject: [PATCH] Added fftplan cache --- modules/core/src/dxt.cpp | 330 +++++++++++++++++------------ modules/core/src/opencl/fft.cl | 55 ++--- modules/core/test/ocl/test_dft.cpp | 4 +- 3 files changed, 225 insertions(+), 164 deletions(-) diff --git a/modules/core/src/dxt.cpp b/modules/core/src/dxt.cpp index de17f07b23..c11b699503 100644 --- a/modules/core/src/dxt.cpp +++ b/modules/core/src/dxt.cpp @@ -2034,50 +2034,6 @@ namespace cv #ifdef HAVE_OPENCL -static bool ocl_packToCCS(InputArray _buffer, OutputArray _dst, int flags) -{ - UMat buffer = _buffer.getUMat(); - UMat dst = _dst.getUMat(); - - buffer = buffer.reshape(1); - if ((flags & DFT_ROWS) == 0 && buffer.rows > 1) - { - // pack to CCS by rows - if (dst.cols > 2) - buffer.colRange(2, dst.cols + (dst.cols % 2)).copyTo(dst.colRange(1, dst.cols-1 + (dst.cols % 2))); - - Mat dst_mat = dst.getMat(ACCESS_WRITE); - Mat buffer_mat = buffer.getMat(ACCESS_READ); - - dst_mat.at(0,0) = buffer_mat.at(0,0); - dst_mat.at(dst_mat.rows-1,0) = buffer_mat.at(buffer.rows/2,0); - for (int i=1; i(i,0) = buffer_mat.at((i+1)/2,0); - dst_mat.at(i+1,0) = buffer_mat.at((i+1)/2,1); - } - - if (dst_mat.cols % 2 == 0) - { - dst_mat.at(0,dst_mat.cols-1) = buffer_mat.at(0,buffer.cols/2); - dst_mat.at(dst_mat.rows-1,dst_mat.cols-1) = buffer_mat.at(buffer.rows/2,buffer.cols/2); - - for (int i=1; i(i,dst_mat.cols-1) = buffer_mat.at((i+1)/2,buffer.cols/2); - dst_mat.at(i+1,dst_mat.cols-1) = buffer_mat.at((i+1)/2,buffer.cols/2+1); - } - } - } - else - { - // pack to CCS each row - buffer.colRange(0,1).copyTo(dst.colRange(0,1)); - buffer.colRange(2, (dst.cols+1)).copyTo(dst.colRange(1, dst.cols)); - } - return true; -} - static std::vector ocl_getRadixes(int cols, int& min_radix) { int factors[34]; @@ -2116,72 +2072,175 @@ static std::vector ocl_getRadixes(int cols, int& min_radix) return radixes; } -static bool ocl_dft_C2C_row(InputArray _src, OutputArray _dst, InputOutputArray _twiddles, int nonzero_rows, int flags) +struct OCL_FftPlan { - int type = _src.type(), depth = CV_MAT_DEPTH(type), channels = CV_MAT_CN(type); - UMat src = _src.getUMat(); + UMat twiddles; + String buildOptions; + int thread_count; - bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; - if (depth == CV_64F && !doubleSupport) - return false; - - int min_radix = INT_MAX; - std::vector radixes = ocl_getRadixes(src.cols, min_radix); - - // generate string with radix calls - String radix_processing; - int n = 1, twiddle_index = 0; - for (size_t i=0; i radixes = ocl_getRadixes(dft_size, min_radix); + thread_count = dft_size / min_radix; + + // generate string with radix calls + String radix_processing; + int n = 1, twiddle_size = 0; + for (size_t i=0; i(); int ptr_index = 0; - int n = 1; + n = 1; for (size_t i=0; idft_size == dft_size) + { + return plan; + } + } + + OCL_FftPlan * newPlan = new OCL_FftPlan(dft_size, flags); + planStorage.push_back(newPlan); + return newPlan; + } + + ~OCL_FftPlanCache() + { + for (std::vector::iterator i = planStorage.begin(), end = planStorage.end(); i != end; ++i) + delete (*i); + planStorage.clear(); + } + +protected: + OCL_FftPlanCache() : + planStorage() + { + } + + std::vector planStorage; +}; + +static bool ocl_packToCCS(InputArray _src, OutputArray _dst, int flags) +{ + UMat src = _src.getUMat(); + _dst.create(src.size(), CV_32F); + UMat dst = _dst.getUMat(); + + src = src.reshape(1); + if ((flags & DFT_ROWS) == 0 && src.rows > 1) + { + // pack to CCS by rows + if (dst.cols > 2) + src.colRange(2, dst.cols + (dst.cols % 2)).copyTo(dst.colRange(1, dst.cols-1 + (dst.cols % 2))); + + Mat dst_mat = dst.getMat(ACCESS_WRITE); + Mat buffer_mat = src.getMat(ACCESS_READ); + + dst_mat.at(0,0) = buffer_mat.at(0,0); + dst_mat.at(dst_mat.rows-1,0) = buffer_mat.at(src.rows/2,0); + for (int i=1; i(i,0) = buffer_mat.at((i+1)/2,0); + dst_mat.at(i+1,0) = buffer_mat.at((i+1)/2,1); + } + + if (dst_mat.cols % 2 == 0) + { + dst_mat.at(0,dst_mat.cols-1) = buffer_mat.at(0,src.cols/2); + dst_mat.at(dst_mat.rows-1,dst_mat.cols-1) = buffer_mat.at(src.rows/2,src.cols/2); + + for (int i=1; i(i,dst_mat.cols-1) = buffer_mat.at((i+1)/2,src.cols/2); + dst_mat.at(i+1,dst_mat.cols-1) = buffer_mat.at((i+1)/2,src.cols/2+1); + } + } + } + else + { + // pack to CCS each row + src.colRange(0,1).copyTo(dst.colRange(0,1)); + src.colRange(2, (dst.cols+1)).copyTo(dst.colRange(1, dst.cols)); + } + return true; +} + +static bool ocl_dft_C2C_row(InputArray _src, OutputArray _dst, int nonzero_rows, int flags) +{ + int type = _src.type(), depth = CV_MAT_DEPTH(type), channels = CV_MAT_CN(type); + + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + if (depth == CV_64F && !doubleSupport) return false; - - k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), ocl::KernelArg::ReadOnlyNoSize(twiddles), thread_count, nonzero_rows); - return k.run(2, globalsize, localsize, false); + + const OCL_FftPlan* plan = OCL_FftPlanCache::getInstance().getFftPlan(_src.cols(), flags); + return plan->enqueueTransform(_src, _dst, nonzero_rows); } static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_rows) @@ -2217,76 +2276,71 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro } } - if (complex_output) + UMat input, output; + if (complex_input) { - //if (is1d) - // _dst.create(Size(src.cols/2+1, src.rows), CV_MAKE_TYPE(depth, 2)); - //else - _dst.create(src.size(), CV_MAKE_TYPE(depth, 2)); + input = src; } else - _dst.create(src.size(), CV_MAKE_TYPE(depth, 1)); + { + if (!inv) + { + // in case real input convert it to complex + input.create(src.size(), CV_MAKE_TYPE(depth, 2)); + std::vector planes; + planes.push_back(src); + planes.push_back(UMat::zeros(src.size(), CV_32F)); + merge(planes, input); + } + else + { + // TODO: unpack from CCS format + } + } + + UMat dst = _dst.getUMat(); - - bool inplace = src.u == dst.u; - //UMat buffer; - - //if (complex_input) - //{ - // if (inplace) - // buffer = src; - // else - // src.copyTo(buffer); - //} - //else - //{ - // if (!inv) - // { - // // in case real input convert it to complex - // buffer.create(src.size(), CV_MAKE_TYPE(depth, 2)); - // std::vector planes; - // planes.push_back(src); - // planes.push_back(UMat::zeros(src.size(), CV_32F)); - // merge(planes, buffer); - // } - // else - // { - // // TODO: unpack from CCS format - // } - //} + if (complex_output) + { + if (real_input && is1d && !inv) + output.create(src.size(), CV_32FC2); + else + output = dst; + } else + { + output.create(src.size(), CV_32FC2); + } if( nonzero_rows <= 0 || nonzero_rows > _src.rows() ) nonzero_rows = _src.rows(); - UMat buffer; - - if (!ocl_dft_C2C_row(src, dst, buffer, nonzero_rows, flags)) + if (!ocl_dft_C2C_row(input, output, nonzero_rows, flags)) return false; if ((flags & DFT_ROWS) == 0 && nonzero_rows > 1) { - transpose(dst, dst); - if (!ocl_dft_C2C_row(dst, dst, buffer, dst.rows, flags)) + transpose(output, output); + if (!ocl_dft_C2C_row(output, output, output.rows, flags)) return false; - transpose(dst, dst); + transpose(output, output); } if (complex_output) { - if (real_input && is1d) - _dst.assign(dst.colRange(0, dst.cols/2+1)); + if (real_input && is1d && !inv) + _dst.assign(output.colRange(0, output.cols/2+1)); else - _dst.assign(dst); + _dst.assign(output); + } + else + { + if (!inv) + ocl_packToCCS(output, _dst, flags); + else + { + // copy real part to dst + } } - //else - //{ - // if (!inv) - // ocl_packToCCS(buffer, _dst, flags); - // else - // { - // // copy real part to dst - // } - //} return true; } diff --git a/modules/core/src/opencl/fft.cl b/modules/core/src/opencl/fft.cl index bd2b863c6c..34da79fafb 100644 --- a/modules/core/src/opencl/fft.cl +++ b/modules/core/src/opencl/fft.cl @@ -28,7 +28,7 @@ float2 twiddle(float2 a) { } __attribute__((always_inline)) -void fft_radix2(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) +void fft_radix2(__local float2* smem, __constant const float2* twiddles, const int x, const int block_size, const int t) { const int k = x & (block_size - 1); float2 a0, a1; @@ -53,17 +53,18 @@ void fft_radix2(__local float2* smem, __global const float2* twiddles, const int } __attribute__((always_inline)) -void fft_radix4(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) +void fft_radix4(__local float2* smem, __constant const float2* twiddles, const int x, const int block_size, const int t) { const int k = x & (block_size - 1); float2 a0, a1, a2, a3; if (x < t) { + const int twiddle_block = block_size / 4; a0 = smem[x]; - a1 = mul_float2(twiddles[3*k],smem[x+t]); - a2 = mul_float2(twiddles[3*k + 1],smem[x+2*t]); - a3 = mul_float2(twiddles[3*k + 2],smem[x+3*t]); + a1 = mul_float2(twiddles[k],smem[x+t]); + a2 = mul_float2(twiddles[k + block_size],smem[x+2*t]); + a3 = mul_float2(twiddles[k + 2*block_size],smem[x+3*t]); } barrier(CLK_LOCAL_MEM_FENCE); @@ -87,7 +88,7 @@ void fft_radix4(__local float2* smem, __global const float2* twiddles, const int } __attribute__((always_inline)) -void fft_radix8(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) +void fft_radix8(__local float2* smem, __constant const float2* twiddles, const int x, const int block_size, const int t) { const int k = x % block_size; float2 a0, a1, a2, a3, a4, a5, a6, a7; @@ -97,13 +98,13 @@ void fft_radix8(__local float2* smem, __global const float2* twiddles, const int int tw_ind = block_size / 8; a0 = smem[x]; - a1 = mul_float2(twiddles[7*k], smem[x + t]); - a2 = mul_float2(twiddles[7*k+1],smem[x+2*t]); - a3 = mul_float2(twiddles[7*k+2],smem[x+3*t]); - a4 = mul_float2(twiddles[7*k+3],smem[x+4*t]); - a5 = mul_float2(twiddles[7*k+4],smem[x+5*t]); - a6 = mul_float2(twiddles[7*k+5],smem[x+6*t]); - a7 = mul_float2(twiddles[7*k+6],smem[x+7*t]); + a1 = mul_float2(twiddles[k], smem[x + t]); + a2 = mul_float2(twiddles[k + block_size],smem[x+2*t]); + a3 = mul_float2(twiddles[k+2*block_size],smem[x+3*t]); + a4 = mul_float2(twiddles[k+3*block_size],smem[x+4*t]); + a5 = mul_float2(twiddles[k+4*block_size],smem[x+5*t]); + a6 = mul_float2(twiddles[k+5*block_size],smem[x+6*t]); + a7 = mul_float2(twiddles[k+6*block_size],smem[x+7*t]); float2 b0, b1, b6, b7; @@ -150,16 +151,23 @@ void fft_radix8(__local float2* smem, __global const float2* twiddles, const int } __attribute__((always_inline)) -void fft_radix3(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) +void fft_radix3(__local float2* smem, __constant const float2* twiddles, const int x, const int block_size, const int t) { const int k = x % block_size; float2 a0, a1, a2; if (x < t) { + //const int twiddle_block = block_size / 3; + //const float theta = -PI * k * 2 / (3 * block_size); + //float2 tw = sincos_float2(theta); + //printf("radix3 %d (%f,%f)(%f,%f)\n", k, tw.x, tw.y, twiddles[k].x, twiddles[k].y); + //tw = sincos_float2(2*theta); + //printf("radix3- %d %d (%f,%f)(%f,%f)\n", k, twiddle_block, tw.x, tw.y, twiddles[k+block_size].x, twiddles[k+block_size].y); + a0 = smem[x]; - a1 = mul_float2(twiddles[2*k], smem[x+t]); - a2 = mul_float2(twiddles[2*k+1], smem[x+2*t]); + a1 = mul_float2(twiddles[k], smem[x+t]); + a2 = mul_float2(twiddles[k+block_size], smem[x+2*t]); } barrier(CLK_LOCAL_MEM_FENCE); @@ -181,7 +189,7 @@ void fft_radix3(__local float2* smem, __global const float2* twiddles, const int } __attribute__((always_inline)) -void fft_radix5(__local float2* smem, __global const float2* twiddles, const int x, const int block_size, const int t) +void fft_radix5(__local float2* smem, __constant const float2* twiddles, const int x, const int block_size, const int t) { const int k = x % block_size; float2 a0, a1, a2, a3, a4; @@ -191,10 +199,10 @@ void fft_radix5(__local float2* smem, __global const float2* twiddles, const int int tw_ind = block_size / 5; a0 = smem[x]; - a1 = mul_float2(twiddles[4*k], smem[x + t]); - a2 = mul_float2(twiddles[4*k+1],smem[x+2*t]); - a3 = mul_float2(twiddles[4*k+2],smem[x+3*t]); - a4 = mul_float2(twiddles[4*k+3],smem[x+4*t]); + a1 = mul_float2(twiddles[k], smem[x + t]); + a2 = mul_float2(twiddles[k + block_size],smem[x+2*t]); + a3 = mul_float2(twiddles[k+2*block_size],smem[x+3*t]); + a4 = mul_float2(twiddles[k+3*block_size],smem[x+4*t]); } barrier(CLK_LOCAL_MEM_FENCE); @@ -237,8 +245,7 @@ void fft_radix5(__local float2* smem, __global const float2* twiddles, const int __kernel void fft_multi_radix(__global const uchar* src_ptr, int src_step, int src_offset, __global uchar* dst_ptr, int dst_step, int dst_offset, - __global const uchar* twiddles_ptr, int twiddles_step, int twiddles_offset, - const int t, const int nz) + __constant float2 * twiddles_ptr, const int t, const int nz) { const int x = get_global_id(0); const int y = get_group_id(1); @@ -248,7 +255,7 @@ __kernel void fft_multi_radix(__global const uchar* src_ptr, int src_step, int s __local float2 smem[LOCAL_SIZE]; __global const float2* src = (__global const float2*)(src_ptr + mad24(y, src_step, mad24(x, (int)(sizeof(float)*2), src_offset))); __global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset))); - __global const float2* twiddles = (__global float2*) twiddles_ptr; + __constant const float2* twiddles = (__constant float2*) twiddles_ptr; const int block_size = LOCAL_SIZE/kercn; #pragma unroll diff --git a/modules/core/test/ocl/test_dft.cpp b/modules/core/test/ocl/test_dft.cpp index 7a7a98852a..2529e949e0 100644 --- a/modules/core/test/ocl/test_dft.cpp +++ b/modules/core/test/ocl/test_dft.cpp @@ -181,9 +181,9 @@ OCL_TEST_P(MulSpectrums, Mat) OCL_INSTANTIATE_TEST_CASE_P(OCL_ImgProc, MulSpectrums, testing::Combine(Bool(), Bool())); -OCL_INSTANTIATE_TEST_CASE_P(Core, Dft, Combine(Values(cv::Size(2, 3), cv::Size(5, 4), cv::Size(30, 20), +OCL_INSTANTIATE_TEST_CASE_P(Core, Dft, Combine(Values(cv::Size(1920, 1), cv::Size(5, 4), cv::Size(30, 20), cv::Size(512, 1), cv::Size(1024, 1024)), - Values((OCL_FFT_TYPE) C2C/*, (OCL_FFT_TYPE) R2R, (OCL_FFT_TYPE) R2C/*, (OCL_FFT_TYPE) C2R*/), + Values(/*(OCL_FFT_TYPE) C2C, (OCL_FFT_TYPE) R2C,*/ (OCL_FFT_TYPE) R2R/*, (OCL_FFT_TYPE) C2R*/), Bool() // DFT_ROWS ) );