Added license header, using cv::Ptr, small fixes.

This commit is contained in:
Alexander Karsakov 2014-07-25 13:11:35 +04:00
parent 66ac46214d
commit 37d01e2d27
3 changed files with 126 additions and 120 deletions

View File

@ -244,7 +244,7 @@ typedef signed char schar;
/* fundamental constants */ /* fundamental constants */
#define CV_PI 3.1415926535897932384626433832795 #define CV_PI 3.1415926535897932384626433832795
#define CV_TWO_PI 6.283185307179586476925286766559 #define CV_2PI 6.283185307179586476925286766559
#define CV_LOG2 0.69314718055994530941723212145818 #define CV_LOG2 0.69314718055994530941723212145818
/****************************************************************************************\ /****************************************************************************************\

View File

@ -1788,12 +1788,136 @@ namespace cv
enum FftType enum FftType
{ {
R2R = 0, R2R = 0, // real to CCS in case forward transform, CCS to real otherwise
C2R = 1, C2R = 1, // complex to real in case inverse transform
R2C = 2, R2C = 2, // real to complex in case forward transform
C2C = 3 C2C = 3 // complex to complex
}; };
struct OCL_FftPlan
{
private:
UMat twiddles;
String buildOptions;
int thread_count;
bool status;
public:
int dft_size;
OCL_FftPlan(int _size): dft_size(_size), status(true)
{
int min_radix;
std::vector<int> radixes, blocks;
ocl_getRadixes(dft_size, radixes, blocks, min_radix);
thread_count = dft_size / min_radix;
if (thread_count > (int) ocl::Device::getDefault().maxWorkGroupSize())
{
status = false;
return;
}
// generate string with radix calls
String radix_processing;
int n = 1, twiddle_size = 0;
for (size_t i=0; i<radixes.size(); i++)
{
int radix = radixes[i], block = blocks[i];
if (block > 1)
radix_processing += format("fft_radix%d_B%d(smem,twiddles+%d,ind,%d,%d);", radix, block, twiddle_size, n, dft_size/radix);
else
radix_processing += format("fft_radix%d(smem,twiddles+%d,ind,%d,%d);", radix, twiddle_size, n, dft_size/radix);
twiddle_size += (radix-1)*n;
n *= radix;
}
Mat tw(1, twiddle_size, CV_32FC2);
float* ptr = tw.ptr<float>();
int ptr_index = 0;
n = 1;
for (size_t i=0; i<radixes.size(); i++)
{
int radix = radixes[i];
n *= radix;
for (int j=1; j<radix; j++)
{
double theta = -CV_2PI*j/n;
for (int k=0; k<(n/radix); k++)
{
ptr[ptr_index++] = (float) cos(k*theta);
ptr[ptr_index++] = (float) sin(k*theta);
}
}
}
twiddles = tw.getUMat(ACCESS_READ);
buildOptions = format("-D LOCAL_SIZE=%d -D kercn=%d -D RADIX_PROCESS=%s",
dft_size, min_radix, radix_processing.c_str());
}
bool enqueueTransform(InputArray _src, OutputArray _dst, int num_dfts, int flags, int fftType, bool rows = true) const
{
if (!status)
return false;
UMat src = _src.getUMat();
UMat dst = _dst.getUMat();
size_t globalsize[2];
size_t localsize[2];
String kernel_name;
bool is1d = (flags & DFT_ROWS) != 0 || num_dfts == 1;
bool inv = (flags & DFT_INVERSE) != 0;
String options = buildOptions;
if (rows)
{
globalsize[0] = thread_count; globalsize[1] = src.rows;
localsize[0] = thread_count; localsize[1] = 1;
kernel_name = !inv ? "fft_multi_radix_rows" : "ifft_multi_radix_rows";
if ((is1d || inv) && (flags & DFT_SCALE))
options += " -D DFT_SCALE";
}
else
{
globalsize[0] = num_dfts; globalsize[1] = thread_count;
localsize[0] = 1; localsize[1] = thread_count;
kernel_name = !inv ? "fft_multi_radix_cols" : "ifft_multi_radix_cols";
if (flags & DFT_SCALE)
options += " -D DFT_SCALE";
}
options += src.channels() == 1 ? " -D REAL_INPUT" : " -D COMPLEX_INPUT";
options += dst.channels() == 1 ? " -D REAL_OUTPUT" : " -D COMPLEX_OUTPUT";
options += is1d ? " -D IS_1D" : "";
if (!inv)
{
if ((is1d && src.channels() == 1) || (rows && (fftType == R2R)))
options += " -D NO_CONJUGATE";
}
else
{
if (rows && (fftType == C2R || fftType == R2R))
options += " -D NO_CONJUGATE";
if (dst.cols % 2 == 0)
options += " -D EVEN";
}
ocl::Kernel k(kernel_name.c_str(), ocl::core::fft_oclsrc, options);
if (k.empty())
return false;
k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrReadOnly(twiddles), thread_count, num_dfts);
return k.run(2, globalsize, localsize, false);
}
private:
static void ocl_getRadixes(int cols, std::vector<int>& radixes, std::vector<int>& blocks, int& min_radix) static void ocl_getRadixes(int cols, std::vector<int>& radixes, std::vector<int>& blocks, int& min_radix)
{ {
int factors[34]; int factors[34];
@ -1862,126 +1986,6 @@ static void ocl_getRadixes(int cols, std::vector<int>& radixes, std::vector<int>
min_radix = min(min_radix, block*radix); min_radix = min(min_radix, block*radix);
} }
} }
struct OCL_FftPlan
{
UMat twiddles;
String buildOptions;
int thread_count;
int dft_size;
bool status;
OCL_FftPlan(int _size): dft_size(_size), status(true)
{
int min_radix;
std::vector<int> radixes, blocks;
ocl_getRadixes(dft_size, radixes, blocks, min_radix);
thread_count = dft_size / min_radix;
if (thread_count > (int) ocl::Device::getDefault().maxWorkGroupSize())
{
status = false;
return;
}
// generate string with radix calls
String radix_processing;
int n = 1, twiddle_size = 0;
for (size_t i=0; i<radixes.size(); i++)
{
int radix = radixes[i], block = blocks[i];
if (block > 1)
radix_processing += format("fft_radix%d_B%d(smem,twiddles+%d,ind,%d,%d);", radix, block, twiddle_size, n, dft_size/radix);
else
radix_processing += format("fft_radix%d(smem,twiddles+%d,ind,%d,%d);", radix, twiddle_size, n, dft_size/radix);
twiddle_size += (radix-1)*n;
n *= radix;
}
Mat tw(1, twiddle_size, CV_32FC2);
float* ptr = tw.ptr<float>();
int ptr_index = 0;
n = 1;
for (size_t i=0; i<radixes.size(); i++)
{
int radix = radixes[i];
n *= radix;
for (int j=1; j<radix; j++)
{
double theta = -CV_TWO_PI*j/n;
for (int k=0; k<(n/radix); k++)
{
ptr[ptr_index++] = (float) cos(k*theta);
ptr[ptr_index++] = (float) sin(k*theta);
}
}
}
twiddles = tw.getUMat(ACCESS_READ);
buildOptions = format("-D LOCAL_SIZE=%d -D kercn=%d -D RADIX_PROCESS=%s",
dft_size, dft_size/thread_count, radix_processing.c_str());
}
bool enqueueTransform(InputArray _src, OutputArray _dst, int num_dfts, int flags, int fftType, bool rows = true) const
{
if (!status)
return false;
UMat src = _src.getUMat();
UMat dst = _dst.getUMat();
size_t globalsize[2];
size_t localsize[2];
String kernel_name;
bool is1d = (flags & DFT_ROWS) != 0 || num_dfts == 1;
bool inv = (flags & DFT_INVERSE) != 0;
String options = buildOptions;
if (rows)
{
globalsize[0] = thread_count; globalsize[1] = src.rows;
localsize[0] = thread_count; localsize[1] = 1;
kernel_name = !inv ? "fft_multi_radix_rows" : "ifft_multi_radix_rows";
if ((is1d || inv) && (flags & DFT_SCALE))
options += " -D DFT_SCALE";
}
else
{
globalsize[0] = num_dfts; globalsize[1] = thread_count;
localsize[0] = 1; localsize[1] = thread_count;
kernel_name = !inv ? "fft_multi_radix_cols" : "ifft_multi_radix_cols";
if (flags & DFT_SCALE)
options += " -D DFT_SCALE";
}
options += src.channels() == 1 ? " -D REAL_INPUT" : " -D COMPLEX_INPUT";
options += dst.channels() == 1 ? " -D REAL_OUTPUT" : " -D COMPLEX_OUTPUT";
options += is1d ? " -D IS_1D" : "";
if (!inv)
{
if ((is1d && src.channels() == 1) || (rows && (fftType == R2R)))
options += " -D NO_CONJUGATE";
}
else
{
if (rows && (fftType == C2R || fftType == R2R))
options += " -D NO_CONJUGATE";
if (dst.cols % 2 == 0)
options += " -D EVEN";
}
ocl::Kernel k(kernel_name.c_str(), ocl::core::fft_oclsrc, options);
if (k.empty())
return false;
k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst), ocl::KernelArg::PtrReadOnly(twiddles), thread_count, num_dfts);
return k.run(2, globalsize, localsize, false);
}
}; };
class OCL_FftPlanCache class OCL_FftPlanCache
@ -1993,27 +1997,24 @@ public:
return planCache; return planCache;
} }
OCL_FftPlan* getFftPlan(int dft_size) Ptr<OCL_FftPlan> getFftPlan(int dft_size)
{ {
for (size_t i = 0, size = planStorage.size(); i < size; ++i) for (size_t i = 0, size = planStorage.size(); i < size; ++i)
{ {
OCL_FftPlan * const plan = planStorage[i]; Ptr<OCL_FftPlan> plan = planStorage[i];
if (plan->dft_size == dft_size) if (plan->dft_size == dft_size)
{ {
return plan; return plan;
} }
} }
OCL_FftPlan * newPlan = new OCL_FftPlan(dft_size); Ptr<OCL_FftPlan> newPlan = Ptr<OCL_FftPlan>(new OCL_FftPlan(dft_size));
planStorage.push_back(newPlan); planStorage.push_back(newPlan);
return newPlan; return newPlan;
} }
~OCL_FftPlanCache() ~OCL_FftPlanCache()
{ {
for (std::vector<OCL_FftPlan *>::iterator i = planStorage.begin(), end = planStorage.end(); i != end; ++i)
delete (*i);
planStorage.clear(); planStorage.clear();
} }
@ -2023,18 +2024,18 @@ protected:
{ {
} }
std::vector<OCL_FftPlan*> planStorage; std::vector<Ptr<OCL_FftPlan> > planStorage;
}; };
static bool ocl_dft_C2C_rows(InputArray _src, OutputArray _dst, int nonzero_rows, int flags, int fftType) static bool ocl_dft_rows(InputArray _src, OutputArray _dst, int nonzero_rows, int flags, int fftType)
{ {
const OCL_FftPlan* plan = OCL_FftPlanCache::getInstance().getFftPlan(_src.cols()); Ptr<OCL_FftPlan> plan = OCL_FftPlanCache::getInstance().getFftPlan(_src.cols());
return plan->enqueueTransform(_src, _dst, nonzero_rows, flags, fftType, true); return plan->enqueueTransform(_src, _dst, nonzero_rows, flags, fftType, true);
} }
static bool ocl_dft_C2C_cols(InputArray _src, OutputArray _dst, int nonzero_cols, int flags, int fftType) static bool ocl_dft_cols(InputArray _src, OutputArray _dst, int nonzero_cols, int flags, int fftType)
{ {
const OCL_FftPlan* plan = OCL_FftPlanCache::getInstance().getFftPlan(_src.rows()); Ptr<OCL_FftPlan> plan = OCL_FftPlanCache::getInstance().getFftPlan(_src.rows());
return plan->enqueueTransform(_src, _dst, nonzero_cols, flags, fftType, false); return plan->enqueueTransform(_src, _dst, nonzero_cols, flags, fftType, false);
} }
@ -2103,13 +2104,13 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
if (!inv) if (!inv)
{ {
if (!ocl_dft_C2C_rows(src, output, nonzero_rows, flags, fftType)) if (!ocl_dft_rows(src, output, nonzero_rows, flags, fftType))
return false; return false;
if (!is1d) if (!is1d)
{ {
int nonzero_cols = fftType == R2R ? output.cols/2 + 1 : output.cols; int nonzero_cols = fftType == R2R ? output.cols/2 + 1 : output.cols;
if (!ocl_dft_C2C_cols(output, _dst, nonzero_cols, flags, fftType)) if (!ocl_dft_cols(output, _dst, nonzero_cols, flags, fftType))
return false; return false;
} }
} }
@ -2118,12 +2119,12 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
if (fftType == C2C) if (fftType == C2C)
{ {
// complex output // complex output
if (!ocl_dft_C2C_rows(src, output, nonzero_rows, flags, fftType)) if (!ocl_dft_rows(src, output, nonzero_rows, flags, fftType))
return false; return false;
if (!is1d) if (!is1d)
{ {
if (!ocl_dft_C2C_cols(output, output, output.cols, flags, fftType)) if (!ocl_dft_cols(output, output, output.cols, flags, fftType))
return false; return false;
} }
} }
@ -2131,16 +2132,16 @@ static bool ocl_dft(InputArray _src, OutputArray _dst, int flags, int nonzero_ro
{ {
if (is1d) if (is1d)
{ {
if (!ocl_dft_C2C_rows(src, output, nonzero_rows, flags, fftType)) if (!ocl_dft_rows(src, output, nonzero_rows, flags, fftType))
return false; return false;
} }
else else
{ {
int nonzero_cols = src.cols/2 + 1; int nonzero_cols = src.cols/2 + 1;
if (!ocl_dft_C2C_cols(src, output, nonzero_cols, flags, fftType)) if (!ocl_dft_cols(src, output, nonzero_cols, flags, fftType))
return false; return false;
if (!ocl_dft_C2C_rows(output, _dst, nonzero_rows, flags, fftType)) if (!ocl_dft_rows(output, _dst, nonzero_rows, flags, fftType))
return false; return false;
} }
} }
@ -2286,7 +2287,7 @@ public:
} }
// no baked plan is found, so let's create a new one // no baked plan is found, so let's create a new one
FftPlan * newPlan = new FftPlan(dft_size, src_step, dst_step, doubleFP, inplace, flags, fftType); Ptr<FftPlan> newPlan = Ptr<FftPlan>(new FftPlan(dft_size, src_step, dst_step, doubleFP, inplace, flags, fftType));
planStorage.push_back(newPlan); planStorage.push_back(newPlan);
return newPlan->plHandle; return newPlan->plHandle;
@ -2294,8 +2295,6 @@ public:
~PlanCache() ~PlanCache()
{ {
for (std::vector<FftPlan *>::iterator i = planStorage.begin(), end = planStorage.end(); i != end; ++i)
delete (*i);
planStorage.clear(); planStorage.clear();
} }
@ -2305,7 +2304,7 @@ protected:
{ {
} }
std::vector<FftPlan *> planStorage; std::vector<Ptr<FftPlan> > planStorage;
}; };
extern "C" { extern "C" {

View File

@ -1,3 +1,10 @@
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.
// Copyright (C) 2014, Itseez, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
#define SQRT_2 0.707106781188f #define SQRT_2 0.707106781188f
#define sin_120 0.866025403784f #define sin_120 0.866025403784f
#define fft5_2 0.559016994374f #define fft5_2 0.559016994374f
@ -509,9 +516,9 @@ void fft_radix5_B2(__local float2* smem, __global const float2* twiddles, const
} }
#ifdef DFT_SCALE #ifdef DFT_SCALE
#define VAL(x, scale) x*scale #define SCALE_VAL(x, scale) x*scale
#else #else
#define VAL(x, scale) x #define SCALE_VAL(x, scale) x
#endif #endif
__kernel void fft_multi_radix_rows(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols, __kernel void fft_multi_radix_rows(__global const uchar* src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
@ -558,15 +565,15 @@ __kernel void fft_multi_radix_rows(__global const uchar* src_ptr, int src_step,
__global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, dst_offset)); __global float2* dst = (__global float2*)(dst_ptr + mad24(y, dst_step, dst_offset));
#pragma unroll #pragma unroll
for (int i=x; i<cols; i+=block_size) for (int i=x; i<cols; i+=block_size)
dst[i] = VAL(smem[i], scale); dst[i] = SCALE_VAL(smem[i], scale);
#else #else
// pack row to CCS // pack row to CCS
__local float* smem_1cn = (__local float*) smem; __local float* smem_1cn = (__local float*) smem;
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, dst_offset)); __global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, dst_offset));
for (int i=x; i<dst_cols-1; i+=block_size) for (int i=x; i<dst_cols-1; i+=block_size)
dst[i+1] = VAL(smem_1cn[i+2], scale); dst[i+1] = SCALE_VAL(smem_1cn[i+2], scale);
if (x == 0) if (x == 0)
dst[0] = VAL(smem_1cn[0], scale); dst[0] = SCALE_VAL(smem_1cn[0], scale);
#endif #endif
} }
else else
@ -611,7 +618,7 @@ __kernel void fft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
__global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset)); __global uchar* dst = dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)*2), dst_offset));
#pragma unroll #pragma unroll
for (int i=0; i<kercn; i++) for (int i=0; i<kercn; i++)
*((__global float2*)(dst + i*block_size*dst_step)) = VAL(smem[y + i*block_size], scale); *((__global float2*)(dst + i*block_size*dst_step)) = SCALE_VAL(smem[y + i*block_size], scale);
#else #else
if (x == 0) if (x == 0)
{ {
@ -619,9 +626,9 @@ __kernel void fft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
__local float* smem_1cn = (__local float*) smem; __local float* smem_1cn = (__local float*) smem;
__global uchar* dst = dst_ptr + mad24(y+1, dst_step, dst_offset); __global uchar* dst = dst_ptr + mad24(y+1, dst_step, dst_offset);
for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size) for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size)
*((__global float*) dst) = VAL(smem_1cn[i+2], scale); *((__global float*) dst) = SCALE_VAL(smem_1cn[i+2], scale);
if (y == 0) if (y == 0)
*((__global float*) (dst_ptr + dst_offset)) = VAL(smem_1cn[0], scale); *((__global float*) (dst_ptr + dst_offset)) = SCALE_VAL(smem_1cn[0], scale);
} }
else if (x == (dst_cols+1)/2) else if (x == (dst_cols+1)/2)
{ {
@ -629,16 +636,16 @@ __kernel void fft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
__local float* smem_1cn = (__local float*) smem; __local float* smem_1cn = (__local float*) smem;
__global uchar* dst = dst_ptr + mad24(dst_cols-1, (int)sizeof(float), mad24(y+1, dst_step, dst_offset)); __global uchar* dst = dst_ptr + mad24(dst_cols-1, (int)sizeof(float), mad24(y+1, dst_step, dst_offset));
for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size) for (int i=y; i<dst_rows-1; i+=block_size, dst+=dst_step*block_size)
*((__global float*) dst) = VAL(smem_1cn[i+2], scale); *((__global float*) dst) = SCALE_VAL(smem_1cn[i+2], scale);
if (y == 0) if (y == 0)
*((__global float*) (dst_ptr + mad24(dst_cols-1, (int)sizeof(float), dst_offset))) = VAL(smem_1cn[0], scale); *((__global float*) (dst_ptr + mad24(dst_cols-1, (int)sizeof(float), dst_offset))) = SCALE_VAL(smem_1cn[0], scale);
} }
else else
{ {
__global uchar* dst = dst_ptr + mad24(x, (int)sizeof(float)*2, mad24(y, dst_step, dst_offset - (int)sizeof(float))); __global uchar* dst = dst_ptr + mad24(x, (int)sizeof(float)*2, mad24(y, dst_step, dst_offset - (int)sizeof(float)));
#pragma unroll #pragma unroll
for (int i=y; i<dst_rows; i+=block_size, dst+=block_size*dst_step) for (int i=y; i<dst_rows; i+=block_size, dst+=block_size*dst_step)
vstore2(VAL(smem[i], scale), 0, (__global float*) dst); vstore2(SCALE_VAL(smem[i], scale), 0, (__global float*) dst);
} }
#endif #endif
} }
@ -724,15 +731,15 @@ __kernel void ifft_multi_radix_rows(__global const uchar* src_ptr, int src_step,
#pragma unroll #pragma unroll
for (int i=0; i<kercn; i++) for (int i=0; i<kercn; i++)
{ {
dst[i*block_size].x = VAL(smem[x + i*block_size].x, scale); dst[i*block_size].x = SCALE_VAL(smem[x + i*block_size].x, scale);
dst[i*block_size].y = VAL(-smem[x + i*block_size].y, scale); dst[i*block_size].y = SCALE_VAL(-smem[x + i*block_size].y, scale);
} }
#else #else
__global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)), dst_offset))); __global float* dst = (__global float*)(dst_ptr + mad24(y, dst_step, mad24(x, (int)(sizeof(float)), dst_offset)));
#pragma unroll #pragma unroll
for (int i=0; i<kercn; i++) for (int i=0; i<kercn; i++)
{ {
dst[i*block_size] = VAL(smem[x + i*block_size].x, scale); dst[i*block_size] = SCALE_VAL(smem[x + i*block_size].x, scale);
} }
#endif #endif
} }
@ -783,9 +790,9 @@ __kernel void ifft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
#pragma unroll #pragma unroll
for (int i=0; i<kercn; i++) for (int i=0; i<kercn; i++)
{ {
__global float2* rez = (__global float2*)(dst + i*block_size*dst_step); __global float2* res = (__global float2*)(dst + i*block_size*dst_step);
rez[0].x = smem[y + i*block_size].x; res[0].x = smem[y + i*block_size].x;
rez[0].y = -smem[y + i*block_size].y; res[0].y = -smem[y + i*block_size].y;
} }
} }
#else #else
@ -848,9 +855,9 @@ __kernel void ifft_multi_radix_cols(__global const uchar* src_ptr, int src_step,
#pragma unroll #pragma unroll
for (int i=0; i<kercn; i++) for (int i=0; i<kercn; i++)
{ {
__global float2* rez = (__global float2*)(dst + i*block_size*dst_step); __global float2* res = (__global float2*)(dst + i*block_size*dst_step);
rez[0].x = smem[y + i*block_size].x; res[0].x = smem[y + i*block_size].x;
rez[0].y = -smem[y + i*block_size].y; res[0].y = -smem[y + i*block_size].y;
} }
} }
#endif #endif