From 29eefe52bb6b927cf6cd2983d429964a4bbef4eb Mon Sep 17 00:00:00 2001 From: peng xiao Date: Tue, 13 Aug 2013 16:27:12 +0800 Subject: [PATCH] Add OpenCL accelerated implementation for Retina. --- modules/bioinspired/CMakeLists.txt | 2 +- .../include/opencv2/bioinspired/retina.hpp | 3 +- modules/bioinspired/src/opencl/retina_kernel.cl | 753 +++++++++ modules/bioinspired/src/precomp.hpp | 6 + modules/bioinspired/src/retina_ocl.cpp | 1651 ++++++++++++++++++++ modules/bioinspired/src/retina_ocl.hpp | 633 ++++++++ modules/bioinspired/test/test_retina_ocl.cpp | 139 ++ 7 files changed, 3185 insertions(+), 2 deletions(-) create mode 100644 modules/bioinspired/src/opencl/retina_kernel.cl create mode 100644 modules/bioinspired/src/retina_ocl.cpp create mode 100644 modules/bioinspired/src/retina_ocl.hpp create mode 100644 modules/bioinspired/test/test_retina_ocl.cpp diff --git a/modules/bioinspired/CMakeLists.txt b/modules/bioinspired/CMakeLists.txt index a27ad73..b0f152c 100644 --- a/modules/bioinspired/CMakeLists.txt +++ b/modules/bioinspired/CMakeLists.txt @@ -1,2 +1,2 @@ set(the_description "Biologically inspired algorithms") -ocv_define_module(bioinspired opencv_core OPTIONAL opencv_highgui) +ocv_define_module(bioinspired opencv_core OPTIONAL opencv_highgui opencv_ocl) diff --git a/modules/bioinspired/include/opencv2/bioinspired/retina.hpp b/modules/bioinspired/include/opencv2/bioinspired/retina.hpp index 4b4c19f..b4fda70 100644 --- a/modules/bioinspired/include/opencv2/bioinspired/retina.hpp +++ b/modules/bioinspired/include/opencv2/bioinspired/retina.hpp @@ -304,7 +304,8 @@ public: CV_EXPORTS Ptr createRetina(Size inputSize); CV_EXPORTS Ptr createRetina(Size inputSize, const bool colorMode, int colorSamplingMethod=RETINA_COLOR_BAYER, const bool useRetinaLogSampling=false, const double reductionFactor=1.0, const double samplingStrenght=10.0); - +CV_EXPORTS Ptr createRetina_OCL(Size inputSize); +CV_EXPORTS Ptr createRetina_OCL(Size inputSize, const bool colorMode, int colorSamplingMethod=RETINA_COLOR_BAYER, const bool useRetinaLogSampling=false, const double reductionFactor=1.0, const double samplingStrenght=10.0); } } #endif /* __OPENCV_BIOINSPIRED_RETINA_HPP__ */ diff --git a/modules/bioinspired/src/opencl/retina_kernel.cl b/modules/bioinspired/src/opencl/retina_kernel.cl new file mode 100644 index 0000000..ed5f831 --- /dev/null +++ b/modules/bioinspired/src/opencl/retina_kernel.cl @@ -0,0 +1,753 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +///////////////////////////////////////////////////////// +//******************************************************* +// basicretinafilter +//////////////// _spatiotemporalLPfilter //////////////// +//_horizontalCausalFilter_addInput +kernel void horizontalCausalFilter_addInput( + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int in_offset, + const int out_offset, + const float _tau, + const float _a +) +{ + int gid = get_global_id(0); + if(gid >= rows) + { + return; + } + + global const float * iptr = + input + mad24(gid, elements_per_row, in_offset / 4); + global float * optr = + output + mad24(gid, elements_per_row, out_offset / 4); + + float res; + float4 in_v4, out_v4, res_v4 = (float4)(0); + //vectorize to increase throughput + for(int i = 0; i < cols / 4; ++i, iptr += 4, optr += 4) + { + in_v4 = vload4(0, iptr); + out_v4 = vload4(0, optr); + + res_v4.x = in_v4.x + _tau * out_v4.x + _a * res_v4.w; + res_v4.y = in_v4.y + _tau * out_v4.y + _a * res_v4.x; + res_v4.z = in_v4.z + _tau * out_v4.z + _a * res_v4.y; + res_v4.w = in_v4.w + _tau * out_v4.w + _a * res_v4.z; + + vstore4(res_v4, 0, optr); + } + res = res_v4.w; + // there may be left some + for(int i = 0; i < cols % 4; ++i, ++iptr, ++optr) + { + res = *iptr + _tau * *optr + _a * res; + *optr = res; + } +} + +//_horizontalAnticausalFilter +kernel void horizontalAnticausalFilter( + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int out_offset, + const float _a +) +{ + int gid = get_global_id(0); + if(gid >= rows) + { + return; + } + + global float * optr = output + + mad24(gid + 1, elements_per_row, - 1 + out_offset / 4); + + float4 result = (float4)(0), out_v4; + // we assume elements_per_row is multple of 4 + for(int i = 0; i < elements_per_row / 4; ++i, optr -= 4) + { + // shift left, `offset` is type `size_t` so it cannot be negative + out_v4 = vload4(0, optr - 3); + + result.w = out_v4.w + _a * result.x; + result.z = out_v4.z + _a * result.w; + result.y = out_v4.y + _a * result.z; + result.x = out_v4.x + _a * result.y; + + vstore4(result, 0, optr - 3); + } +} + +//_verticalCausalFilter +kernel void verticalCausalFilter( + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int out_offset, + const float _a +) +{ + int gid = get_global_id(0); + if(gid >= cols) + { + return; + } + + global float * optr = output + gid + out_offset / 4; + float result = 0; + for(int i = 0; i < rows; ++i, optr += elements_per_row) + { + result = *optr + _a * result; + *optr = result; + } +} + +//_verticalCausalFilter +kernel void verticalAnticausalFilter_multGain( + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int out_offset, + const float _a, + const float _gain +) +{ + int gid = get_global_id(0); + if(gid >= cols) + { + return; + } + + global float * optr = output + (rows - 1) * elements_per_row + gid + out_offset / 4; + float result = 0; + for(int i = 0; i < rows; ++i, optr -= elements_per_row) + { + result = *optr + _a * result; + *optr = _gain * result; + } +} +// +// end of _spatiotemporalLPfilter +///////////////////////////////////////////////////////////////////// + +//////////////// horizontalAnticausalFilter_Irregular //////////////// +kernel void horizontalAnticausalFilter_Irregular( + global float * output, + global float * buffer, + const int cols, + const int rows, + const int elements_per_row, + const int out_offset, + const int buffer_offset +) +{ + int gid = get_global_id(0); + if(gid >= rows) + { + return; + } + + global float * optr = + output + mad24(rows - gid, elements_per_row, -1 + out_offset / 4); + global float * bptr = + buffer + mad24(rows - gid, elements_per_row, -1 + buffer_offset / 4); + + float4 buf_v4, out_v4, res_v4 = (float4)(0); + + for(int i = 0; i < elements_per_row / 4; ++i, optr -= 4, bptr -= 4) + { + buf_v4 = vload4(0, bptr - 3); + out_v4 = vload4(0, optr - 3); + + res_v4.w = out_v4.w + buf_v4.w * res_v4.x; + res_v4.z = out_v4.z + buf_v4.z * res_v4.w; + res_v4.y = out_v4.y + buf_v4.y * res_v4.z; + res_v4.x = out_v4.x + buf_v4.x * res_v4.y; + + vstore4(res_v4, 0, optr - 3); + } +} + +//////////////// verticalCausalFilter_Irregular //////////////// +kernel void verticalCausalFilter_Irregular( + global float * output, + global float * buffer, + const int cols, + const int rows, + const int elements_per_row, + const int out_offset, + const int buffer_offset +) +{ + int gid = get_global_id(0); + if(gid >= cols) + { + return; + } + + global float * optr = output + gid + out_offset / 4; + global float * bptr = buffer + gid + buffer_offset / 4; + float result = 0; + for(int i = 0; i < rows; ++i, optr += elements_per_row, bptr += elements_per_row) + { + result = *optr + *bptr * result; + *optr = result; + } +} + +//////////////// _adaptiveHorizontalCausalFilter_addInput //////////////// +kernel void adaptiveHorizontalCausalFilter_addInput( + global const float * input, + global const float * gradient, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int in_offset, + const int grad_offset, + const int out_offset +) +{ + int gid = get_global_id(0); + if(gid >= rows) + { + return; + } + + global const float * iptr = + input + mad24(gid, elements_per_row, in_offset / 4); + global const float * gptr = + gradient + mad24(gid, elements_per_row, grad_offset / 4); + global float * optr = + output + mad24(gid, elements_per_row, out_offset / 4); + + float4 in_v4, grad_v4, out_v4, res_v4 = (float4)(0); + for(int i = 0; i < cols / 4; ++i, iptr += 4, gptr += 4, optr += 4) + { + in_v4 = vload4(0, iptr); + grad_v4 = vload4(0, gptr); + + res_v4.x = in_v4.x + grad_v4.x * res_v4.w; + res_v4.y = in_v4.y + grad_v4.y * res_v4.x; + res_v4.z = in_v4.z + grad_v4.z * res_v4.y; + res_v4.w = in_v4.w + grad_v4.w * res_v4.z; + + vstore4(res_v4, 0, optr); + } + for(int i = 0; i < cols % 4; ++i, ++iptr, ++gptr, ++optr) + { + res_v4.w = *iptr + *gptr * res_v4.w; + *optr = res_v4.w; + } +} + +//////////////// _adaptiveVerticalAnticausalFilter_multGain //////////////// +kernel void adaptiveVerticalAnticausalFilter_multGain( + global const float * gradient, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const int grad_offset, + const int out_offset, + const float gain +) +{ + int gid = get_global_id(0); + if(gid >= cols) + { + return; + } + + int start_idx = mad24(rows - 1, elements_per_row, gid); + + global const float * gptr = gradient + start_idx + grad_offset / 4; + global float * optr = output + start_idx + out_offset / 4; + + float result = 0; + for(int i = 0; i < rows; ++i, gptr -= elements_per_row, optr -= elements_per_row) + { + result = *optr + *gptr * result; + *optr = gain * result; + } +} + +//////////////// _localLuminanceAdaptation //////////////// +// FIXME: +// This kernel seems to have precision problem on GPU +kernel void localLuminanceAdaptation( + global const float * luma, + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const float _localLuminanceAddon, + const float _localLuminanceFactor, + const float _maxInputValue +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + int offset = mad24(gidy, elements_per_row, gidx); + + float X0 = luma[offset] * _localLuminanceFactor + _localLuminanceAddon; + float input_val = input[offset]; + // output of the following line may be different between GPU and CPU + output[offset] = (_maxInputValue + X0) * input_val / (input_val + X0 + 0.00000000001f); +} +// end of basicretinafilter +//******************************************************* +///////////////////////////////////////////////////////// + + + +///////////////////////////////////////////////////////// +//****************************************************** +// magno +// TODO: this kernel has too many buffer accesses, better to make it +// vector read/write for fetch efficiency +kernel void amacrineCellsComputing( + global const float * opl_on, + global const float * opl_off, + global float * prev_in_on, + global float * prev_in_off, + global float * out_on, + global float * out_off, + const int cols, + const int rows, + const int elements_per_row, + const float coeff +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + + int offset = mad24(gidy, elements_per_row, gidx); + opl_on += offset; + opl_off += offset; + prev_in_on += offset; + prev_in_off += offset; + out_on += offset; + out_off += offset; + + float magnoXonPixelResult = coeff * (*out_on + *opl_on - *prev_in_on); + *out_on = fmax(magnoXonPixelResult, 0); + float magnoXoffPixelResult = coeff * (*out_off + *opl_off - *prev_in_off); + *out_off = fmax(magnoXoffPixelResult, 0); + + *prev_in_on = *opl_on; + *prev_in_off = *opl_off; +} + +///////////////////////////////////////////////////////// +//****************************************************** +// parvo +// TODO: this kernel has too many buffer accesses, needs optimization +kernel void OPL_OnOffWaysComputing( + global float4 * photo_out, + global float4 * horiz_out, + global float4 * bipol_on, + global float4 * bipol_off, + global float4 * parvo_on, + global float4 * parvo_off, + const int cols, + const int rows, + const int elements_per_row +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx * 4 >= cols || gidy >= rows) + { + return; + } + // we assume elements_per_row must be multiples of 4 + int offset = mad24(gidy, elements_per_row >> 2, gidx); + photo_out += offset; + horiz_out += offset; + bipol_on += offset; + bipol_off += offset; + parvo_on += offset; + parvo_off += offset; + + float4 diff = *photo_out - *horiz_out; + float4 isPositive;// = convert_float4(diff > (float4)(0.0f, 0.0f, 0.0f, 0.0f)); + isPositive.x = diff.x > 0.0f; + isPositive.y = diff.y > 0.0f; + isPositive.z = diff.z > 0.0f; + isPositive.w = diff.w > 0.0f; + float4 res_on = isPositive * diff; + float4 res_off = (isPositive - (float4)(1.0f)) * diff; + + *bipol_on = res_on; + *parvo_on = res_on; + + *bipol_off = res_off; + *parvo_off = res_off; +} + +///////////////////////////////////////////////////////// +//****************************************************** +// retinacolor +inline int bayerSampleOffset(int step, int rows, int x, int y) +{ + return mad24(y, step, x) + + ((y % 2) + (x % 2)) * rows * step; +} + + +/////// colorMultiplexing ////// +kernel void runColorMultiplexingBayer( + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + + int offset = mad24(gidy, elements_per_row, gidx); + output[offset] = input[bayerSampleOffset(elements_per_row, rows, gidx, gidy)]; +} + +kernel void runColorDemultiplexingBayer( + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + + int offset = mad24(gidy, elements_per_row, gidx); + output[bayerSampleOffset(elements_per_row, rows, gidx, gidy)] = input[offset]; +} + +kernel void demultiplexAssign( + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row +) +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + + int offset = bayerSampleOffset(elements_per_row, rows, gidx, gidy); + output[offset] = input[offset]; +} + + +//// normalizeGrayOutputCentredSigmoide +kernel void normalizeGrayOutputCentredSigmoide( + global const float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const float meanval, + const float X0 +) + +{ + int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + int offset = mad24(gidy, elements_per_row, gidx); + + float input_val = input[offset]; + output[offset] = meanval + + (meanval + X0) * (input_val - meanval) / (fabs(input_val - meanval) + X0); +} + +//// normalize by photoreceptors density +kernel void normalizePhotoDensity( + global const float * chroma, + global const float * colorDensity, + global const float * multiplex, + global float * luma, + global float * demultiplex, + const int cols, + const int rows, + const int elements_per_row, + const float pG +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + int index = offset; + + float Cr = chroma[index] * colorDensity[index]; + index += elements_per_row * rows; + float Cg = chroma[index] * colorDensity[index]; + index += elements_per_row * rows; + float Cb = chroma[index] * colorDensity[index]; + + const float luma_res = (Cr + Cg + Cb) * pG; + luma[offset] = luma_res; + demultiplex[bayerSampleOffset(elements_per_row, rows, gidx, gidy)] = + multiplex[offset] - luma_res; +} + + + +//////// computeGradient /////// +// TODO: +// this function maybe accelerated by image2d_t or lds +kernel void computeGradient( + global const float * luma, + global float * gradient, + const int cols, + const int rows, + const int elements_per_row +) +{ + int gidx = get_global_id(0) + 2, gidy = get_global_id(1) + 2; + if(gidx >= cols - 2 || gidy >= rows - 2) + { + return; + } + int offset = mad24(gidy, elements_per_row, gidx); + luma += offset; + + // horizontal and vertical local gradients + const float v_grad = fabs(luma[elements_per_row] - luma[- elements_per_row]); + const float h_grad = fabs(luma[1] - luma[-1]); + + // neighborhood horizontal and vertical gradients + const float cur_val = luma[0]; + const float v_grad_p = fabs(cur_val - luma[- 2 * elements_per_row]); + const float h_grad_p = fabs(cur_val - luma[- 2]); + const float v_grad_n = fabs(cur_val - luma[2 * elements_per_row]); + const float h_grad_n = fabs(cur_val - luma[2]); + + const float horiz_grad = 0.5f * h_grad + 0.25f * (h_grad_p + h_grad_n); + const float verti_grad = 0.5f * v_grad + 0.25f * (v_grad_p + v_grad_n); + const bool is_vertical_greater = horiz_grad < verti_grad; + + gradient[offset + elements_per_row * rows] = is_vertical_greater ? 0.06f : 0.57f; + gradient[offset ] = is_vertical_greater ? 0.57f : 0.06f; +} + + +/////// substractResidual /////// +kernel void substractResidual( + global float * input, + const int cols, + const int rows, + const int elements_per_row, + const float pR, + const float pG, + const float pB +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + int indices [3] = + { + mad24(gidy, elements_per_row, gidx), + mad24(gidy + rows, elements_per_row, gidx), + mad24(gidy + 2 * rows, elements_per_row, gidx) + }; + float vals[3] = {input[indices[0]], input[indices[1]], input[indices[2]]}; + float residu = pR * vals[0] + pG * vals[1] + pB * vals[2]; + + input[indices[0]] = vals[0] - residu; + input[indices[1]] = vals[1] - residu; + input[indices[2]] = vals[2] - residu; +} + +///// clipRGBOutput_0_maxInputValue ///// +kernel void clipRGBOutput_0_maxInputValue( + global float * input, + const int cols, + const int rows, + const int elements_per_row, + const float maxVal +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + float val = input[offset]; + val = clamp(val, 0.0f, maxVal); + input[offset] = val; +} + +//// normalizeGrayOutputNearZeroCentreredSigmoide //// +kernel void normalizeGrayOutputNearZeroCentreredSigmoide( + global float * input, + global float * output, + const int cols, + const int rows, + const int elements_per_row, + const float maxVal, + const float X0cube +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + float currentCubeLuminance = input[offset]; + currentCubeLuminance = currentCubeLuminance * currentCubeLuminance * currentCubeLuminance; + output[offset] = currentCubeLuminance * X0cube / (X0cube + currentCubeLuminance); +} + +//// centerReductImageLuminance //// +kernel void centerReductImageLuminance( + global float * input, + const int cols, + const int rows, + const int elements_per_row, + const float mean, + const float std_dev +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + + float val = input[offset]; + input[offset] = (val - mean) / std_dev; +} + +//// inverseValue //// +kernel void inverseValue( + global float * input, + const int cols, + const int rows, + const int elements_per_row +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + input[offset] = 1.f / input[offset]; +} + +#define CV_PI 3.1415926535897932384626433832795 + +//// _processRetinaParvoMagnoMapping //// +kernel void processRetinaParvoMagnoMapping( + global float * parvo, + global float * magno, + global float * output, + const int cols, + const int rows, + const int halfCols, + const int halfRows, + const int elements_per_row, + const float minDistance +) +{ + const int gidx = get_global_id(0), gidy = get_global_id(1); + if(gidx >= cols || gidy >= rows) + { + return; + } + const int offset = mad24(gidy, elements_per_row, gidx); + + float distanceToCenter = + sqrt(((float)(gidy - halfRows) * (gidy - halfRows) + (gidx - halfCols) * (gidx - halfCols))); + + float a = distanceToCenter < minDistance ? + (0.5f + 0.5f * (float)cos(CV_PI * distanceToCenter / minDistance)) : 0; + float b = 1.f - a; + + output[offset] = parvo[offset] * a + magno[offset] * b; +} diff --git a/modules/bioinspired/src/precomp.hpp b/modules/bioinspired/src/precomp.hpp index 066e15b..541b970 100644 --- a/modules/bioinspired/src/precomp.hpp +++ b/modules/bioinspired/src/precomp.hpp @@ -43,11 +43,17 @@ #ifndef __OPENCV_PRECOMP_H__ #define __OPENCV_PRECOMP_H__ +#include "opencv2/opencv_modules.hpp" #include "opencv2/bioinspired.hpp" #include "opencv2/core/utility.hpp" #include "opencv2/core/private.hpp" #include + +#ifdef HAVE_OPENCV_OCL + #include "opencv2/ocl/private/util.hpp" +#endif + namespace cv { diff --git a/modules/bioinspired/src/retina_ocl.cpp b/modules/bioinspired/src/retina_ocl.cpp new file mode 100644 index 0000000..ba98da2 --- /dev/null +++ b/modules/bioinspired/src/retina_ocl.cpp @@ -0,0 +1,1651 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "precomp.hpp" +#include "retina_ocl.hpp" +#include +#include + +#ifdef HAVE_OPENCV_OCL + +#define NOT_IMPLEMENTED CV_Error(cv::Error::StsNotImplemented, "Not implemented") + +namespace cv +{ +namespace ocl +{ +//OpenCL kernel file string pointer +extern const char * retina_kernel; +} +} + +namespace cv +{ +namespace bioinspired +{ +namespace ocl +{ +using namespace cv::ocl; + +class RetinaOCLImpl : public Retina +{ +public: + RetinaOCLImpl(Size getInputSize); + RetinaOCLImpl(Size getInputSize, const bool colorMode, int colorSamplingMethod = RETINA_COLOR_BAYER, const bool useRetinaLogSampling = false, const double reductionFactor = 1.0, const double samplingStrenght = 10.0); + virtual ~RetinaOCLImpl(); + + Size getInputSize(); + Size getOutputSize(); + + void setup(String retinaParameterFile = "", const bool applyDefaultSetupOnFailure = true); + void setup(cv::FileStorage &fs, const bool applyDefaultSetupOnFailure = true); + void setup(RetinaParameters newParameters); + + RetinaOCLImpl::RetinaParameters getParameters(); + + const String printSetup(); + virtual void write( String fs ) const; + virtual void write( FileStorage& fs ) const; + + void setupOPLandIPLParvoChannel(const bool colorMode = true, const bool normaliseOutput = true, const float photoreceptorsLocalAdaptationSensitivity = 0.7, const float photoreceptorsTemporalConstant = 0.5, const float photoreceptorsSpatialConstant = 0.53, const float horizontalCellsGain = 0, const float HcellsTemporalConstant = 1, const float HcellsSpatialConstant = 7, const float ganglionCellsSensitivity = 0.7); + void setupIPLMagnoChannel(const bool normaliseOutput = true, const float parasolCells_beta = 0, const float parasolCells_tau = 0, const float parasolCells_k = 7, const float amacrinCellsTemporalCutFrequency = 1.2, const float V0CompressionParameter = 0.95, const float localAdaptintegration_tau = 0, const float localAdaptintegration_k = 7); + + void run(InputArray inputImage); + void getParvo(OutputArray retinaOutput_parvo); + void getMagno(OutputArray retinaOutput_magno); + + void setColorSaturation(const bool saturateColors = true, const float colorSaturationValue = 4.0); + void clearBuffers(); + void activateMovingContoursProcessing(const bool activate); + void activateContoursProcessing(const bool activate); + + // unimplemented interfaces: + void applyFastToneMapping(InputArray /*inputImage*/, OutputArray /*outputToneMappedImage*/) { NOT_IMPLEMENTED; } + void getParvoRAW(OutputArray /*retinaOutput_parvo*/) { NOT_IMPLEMENTED; } + void getMagnoRAW(OutputArray /*retinaOutput_magno*/) { NOT_IMPLEMENTED; } + const Mat getMagnoRAW() const { NOT_IMPLEMENTED; return Mat(); } + const Mat getParvoRAW() const { NOT_IMPLEMENTED; return Mat(); } + +protected: + RetinaParameters _retinaParameters; + cv::ocl::oclMat _inputBuffer; + RetinaFilter* _retinaFilter; + bool convertToColorPlanes(const cv::ocl::oclMat& input, cv::ocl::oclMat &output); + void convertToInterleaved(const cv::ocl::oclMat& input, bool colorMode, cv::ocl::oclMat &output); + void _init(const Size getInputSize, const bool colorMode, int colorSamplingMethod = RETINA_COLOR_BAYER, const bool useRetinaLogSampling = false, const double reductionFactor = 1.0, const double samplingStrenght = 10.0); +}; + +RetinaOCLImpl::RetinaOCLImpl(const cv::Size inputSz) +{ + _retinaFilter = 0; + _init(inputSz, true, RETINA_COLOR_BAYER, false); +} + +RetinaOCLImpl::RetinaOCLImpl(const cv::Size inputSz, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght) +{ + _retinaFilter = 0; + _init(inputSz, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght); +}; + +RetinaOCLImpl::~RetinaOCLImpl() +{ + if (_retinaFilter) + { + delete _retinaFilter; + } +} + +/** +* retreive retina input buffer size +*/ +Size RetinaOCLImpl::getInputSize() +{ + return cv::Size(_retinaFilter->getInputNBcolumns(), _retinaFilter->getInputNBrows()); +} + +/** +* retreive retina output buffer size +*/ +Size RetinaOCLImpl::getOutputSize() +{ + return cv::Size(_retinaFilter->getOutputNBcolumns(), _retinaFilter->getOutputNBrows()); +} + + +void RetinaOCLImpl::setColorSaturation(const bool saturateColors, const float colorSaturationValue) +{ + _retinaFilter->setColorSaturation(saturateColors, colorSaturationValue); +} + +struct RetinaOCLImpl::RetinaParameters RetinaOCLImpl::getParameters() +{ + return _retinaParameters; +} + + +void RetinaOCLImpl::setup(String retinaParameterFile, const bool applyDefaultSetupOnFailure) +{ + try + { + // opening retinaParameterFile in read mode + cv::FileStorage fs(retinaParameterFile, cv::FileStorage::READ); + setup(fs, applyDefaultSetupOnFailure); + } + catch(Exception &e) + { + std::cout << "RetinaOCLImpl::setup: wrong/unappropriate xml parameter file : error report :`n=>" << e.what() << std::endl; + if (applyDefaultSetupOnFailure) + { + std::cout << "RetinaOCLImpl::setup: resetting retina with default parameters" << std::endl; + setupOPLandIPLParvoChannel(); + setupIPLMagnoChannel(); + } + else + { + std::cout << "=> keeping current parameters" << std::endl; + } + } +} + +void RetinaOCLImpl::setup(cv::FileStorage &fs, const bool applyDefaultSetupOnFailure) +{ + try + { + // read parameters file if it exists or apply default setup if asked for + if (!fs.isOpened()) + { + std::cout << "RetinaOCLImpl::setup: provided parameters file could not be open... skeeping configuration" << std::endl; + return; + // implicit else case : retinaParameterFile could be open (it exists at least) + } + // OPL and Parvo init first... update at the same time the parameters structure and the retina core + cv::FileNode rootFn = fs.root(), currFn = rootFn["OPLandIPLparvo"]; + currFn["colorMode"] >> _retinaParameters.OPLandIplParvo.colorMode; + currFn["normaliseOutput"] >> _retinaParameters.OPLandIplParvo.normaliseOutput; + currFn["photoreceptorsLocalAdaptationSensitivity"] >> _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity; + currFn["photoreceptorsTemporalConstant"] >> _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant; + currFn["photoreceptorsSpatialConstant"] >> _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant; + currFn["horizontalCellsGain"] >> _retinaParameters.OPLandIplParvo.horizontalCellsGain; + currFn["hcellsTemporalConstant"] >> _retinaParameters.OPLandIplParvo.hcellsTemporalConstant; + currFn["hcellsSpatialConstant"] >> _retinaParameters.OPLandIplParvo.hcellsSpatialConstant; + currFn["ganglionCellsSensitivity"] >> _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity; + setupOPLandIPLParvoChannel(_retinaParameters.OPLandIplParvo.colorMode, _retinaParameters.OPLandIplParvo.normaliseOutput, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant, _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant, _retinaParameters.OPLandIplParvo.horizontalCellsGain, _retinaParameters.OPLandIplParvo.hcellsTemporalConstant, _retinaParameters.OPLandIplParvo.hcellsSpatialConstant, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity); + + // init retina IPL magno setup... update at the same time the parameters structure and the retina core + currFn = rootFn["IPLmagno"]; + currFn["normaliseOutput"] >> _retinaParameters.IplMagno.normaliseOutput; + currFn["parasolCells_beta"] >> _retinaParameters.IplMagno.parasolCells_beta; + currFn["parasolCells_tau"] >> _retinaParameters.IplMagno.parasolCells_tau; + currFn["parasolCells_k"] >> _retinaParameters.IplMagno.parasolCells_k; + currFn["amacrinCellsTemporalCutFrequency"] >> _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency; + currFn["V0CompressionParameter"] >> _retinaParameters.IplMagno.V0CompressionParameter; + currFn["localAdaptintegration_tau"] >> _retinaParameters.IplMagno.localAdaptintegration_tau; + currFn["localAdaptintegration_k"] >> _retinaParameters.IplMagno.localAdaptintegration_k; + + setupIPLMagnoChannel(_retinaParameters.IplMagno.normaliseOutput, _retinaParameters.IplMagno.parasolCells_beta, _retinaParameters.IplMagno.parasolCells_tau, _retinaParameters.IplMagno.parasolCells_k, _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency, _retinaParameters.IplMagno.V0CompressionParameter, _retinaParameters.IplMagno.localAdaptintegration_tau, _retinaParameters.IplMagno.localAdaptintegration_k); + + } + catch(Exception &e) + { + std::cout << "RetinaOCLImpl::setup: resetting retina with default parameters" << std::endl; + if (applyDefaultSetupOnFailure) + { + setupOPLandIPLParvoChannel(); + setupIPLMagnoChannel(); + } + std::cout << "RetinaOCLImpl::setup: wrong/unappropriate xml parameter file : error report :`n=>" << e.what() << std::endl; + std::cout << "=> keeping current parameters" << std::endl; + } +} + +void RetinaOCLImpl::setup(cv::bioinspired::Retina::RetinaParameters newConfiguration) +{ + // simply copy structures + memcpy(&_retinaParameters, &newConfiguration, sizeof(cv::bioinspired::Retina::RetinaParameters)); + // apply setup + setupOPLandIPLParvoChannel(_retinaParameters.OPLandIplParvo.colorMode, _retinaParameters.OPLandIplParvo.normaliseOutput, _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity, _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant, _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant, _retinaParameters.OPLandIplParvo.horizontalCellsGain, _retinaParameters.OPLandIplParvo.hcellsTemporalConstant, _retinaParameters.OPLandIplParvo.hcellsSpatialConstant, _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity); + setupIPLMagnoChannel(_retinaParameters.IplMagno.normaliseOutput, _retinaParameters.IplMagno.parasolCells_beta, _retinaParameters.IplMagno.parasolCells_tau, _retinaParameters.IplMagno.parasolCells_k, _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency, _retinaParameters.IplMagno.V0CompressionParameter, _retinaParameters.IplMagno.localAdaptintegration_tau, _retinaParameters.IplMagno.localAdaptintegration_k); +} + +const String RetinaOCLImpl::printSetup() +{ + std::stringstream outmessage; + + // displaying OPL and IPL parvo setup + outmessage << "Current Retina instance setup :" + << "\nOPLandIPLparvo" << "{" + << "\n==> colorMode : " << _retinaParameters.OPLandIplParvo.colorMode + << "\n==> normalizeParvoOutput :" << _retinaParameters.OPLandIplParvo.normaliseOutput + << "\n==> photoreceptorsLocalAdaptationSensitivity : " << _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity + << "\n==> photoreceptorsTemporalConstant : " << _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant + << "\n==> photoreceptorsSpatialConstant : " << _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant + << "\n==> horizontalCellsGain : " << _retinaParameters.OPLandIplParvo.horizontalCellsGain + << "\n==> hcellsTemporalConstant : " << _retinaParameters.OPLandIplParvo.hcellsTemporalConstant + << "\n==> hcellsSpatialConstant : " << _retinaParameters.OPLandIplParvo.hcellsSpatialConstant + << "\n==> parvoGanglionCellsSensitivity : " << _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity + << "}\n"; + + // displaying IPL magno setup + outmessage << "Current Retina instance setup :" + << "\nIPLmagno" << "{" + << "\n==> normaliseOutput : " << _retinaParameters.IplMagno.normaliseOutput + << "\n==> parasolCells_beta : " << _retinaParameters.IplMagno.parasolCells_beta + << "\n==> parasolCells_tau : " << _retinaParameters.IplMagno.parasolCells_tau + << "\n==> parasolCells_k : " << _retinaParameters.IplMagno.parasolCells_k + << "\n==> amacrinCellsTemporalCutFrequency : " << _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency + << "\n==> V0CompressionParameter : " << _retinaParameters.IplMagno.V0CompressionParameter + << "\n==> localAdaptintegration_tau : " << _retinaParameters.IplMagno.localAdaptintegration_tau + << "\n==> localAdaptintegration_k : " << _retinaParameters.IplMagno.localAdaptintegration_k + << "}"; + return outmessage.str().c_str(); +} + +void RetinaOCLImpl::write( String fs ) const +{ + FileStorage parametersSaveFile(fs, cv::FileStorage::WRITE ); + write(parametersSaveFile); +} + +void RetinaOCLImpl::write( FileStorage& fs ) const +{ + if (!fs.isOpened()) + { + return; // basic error case + } + fs << "OPLandIPLparvo" << "{"; + fs << "colorMode" << _retinaParameters.OPLandIplParvo.colorMode; + fs << "normaliseOutput" << _retinaParameters.OPLandIplParvo.normaliseOutput; + fs << "photoreceptorsLocalAdaptationSensitivity" << _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity; + fs << "photoreceptorsTemporalConstant" << _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant; + fs << "photoreceptorsSpatialConstant" << _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant; + fs << "horizontalCellsGain" << _retinaParameters.OPLandIplParvo.horizontalCellsGain; + fs << "hcellsTemporalConstant" << _retinaParameters.OPLandIplParvo.hcellsTemporalConstant; + fs << "hcellsSpatialConstant" << _retinaParameters.OPLandIplParvo.hcellsSpatialConstant; + fs << "ganglionCellsSensitivity" << _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity; + fs << "}"; + fs << "IPLmagno" << "{"; + fs << "normaliseOutput" << _retinaParameters.IplMagno.normaliseOutput; + fs << "parasolCells_beta" << _retinaParameters.IplMagno.parasolCells_beta; + fs << "parasolCells_tau" << _retinaParameters.IplMagno.parasolCells_tau; + fs << "parasolCells_k" << _retinaParameters.IplMagno.parasolCells_k; + fs << "amacrinCellsTemporalCutFrequency" << _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency; + fs << "V0CompressionParameter" << _retinaParameters.IplMagno.V0CompressionParameter; + fs << "localAdaptintegration_tau" << _retinaParameters.IplMagno.localAdaptintegration_tau; + fs << "localAdaptintegration_k" << _retinaParameters.IplMagno.localAdaptintegration_k; + fs << "}"; +} + +void RetinaOCLImpl::setupOPLandIPLParvoChannel(const bool colorMode, const bool normaliseOutput, const float photoreceptorsLocalAdaptationSensitivity, const float photoreceptorsTemporalConstant, const float photoreceptorsSpatialConstant, const float horizontalCellsGain, const float HcellsTemporalConstant, const float HcellsSpatialConstant, const float ganglionCellsSensitivity) +{ + // retina core parameters setup + _retinaFilter->setColorMode(colorMode); + _retinaFilter->setPhotoreceptorsLocalAdaptationSensitivity(photoreceptorsLocalAdaptationSensitivity); + _retinaFilter->setOPLandParvoParameters(0, photoreceptorsTemporalConstant, photoreceptorsSpatialConstant, horizontalCellsGain, HcellsTemporalConstant, HcellsSpatialConstant, ganglionCellsSensitivity); + _retinaFilter->setParvoGanglionCellsLocalAdaptationSensitivity(ganglionCellsSensitivity); + _retinaFilter->activateNormalizeParvoOutput_0_maxOutputValue(normaliseOutput); + + // update parameters struture + + _retinaParameters.OPLandIplParvo.colorMode = colorMode; + _retinaParameters.OPLandIplParvo.normaliseOutput = normaliseOutput; + _retinaParameters.OPLandIplParvo.photoreceptorsLocalAdaptationSensitivity = photoreceptorsLocalAdaptationSensitivity; + _retinaParameters.OPLandIplParvo.photoreceptorsTemporalConstant = photoreceptorsTemporalConstant; + _retinaParameters.OPLandIplParvo.photoreceptorsSpatialConstant = photoreceptorsSpatialConstant; + _retinaParameters.OPLandIplParvo.horizontalCellsGain = horizontalCellsGain; + _retinaParameters.OPLandIplParvo.hcellsTemporalConstant = HcellsTemporalConstant; + _retinaParameters.OPLandIplParvo.hcellsSpatialConstant = HcellsSpatialConstant; + _retinaParameters.OPLandIplParvo.ganglionCellsSensitivity = ganglionCellsSensitivity; +} + +void RetinaOCLImpl::setupIPLMagnoChannel(const bool normaliseOutput, const float parasolCells_beta, const float parasolCells_tau, const float parasolCells_k, const float amacrinCellsTemporalCutFrequency, const float V0CompressionParameter, const float localAdaptintegration_tau, const float localAdaptintegration_k) +{ + + _retinaFilter->setMagnoCoefficientsTable(parasolCells_beta, parasolCells_tau, parasolCells_k, amacrinCellsTemporalCutFrequency, V0CompressionParameter, localAdaptintegration_tau, localAdaptintegration_k); + _retinaFilter->activateNormalizeMagnoOutput_0_maxOutputValue(normaliseOutput); + + // update parameters struture + _retinaParameters.IplMagno.normaliseOutput = normaliseOutput; + _retinaParameters.IplMagno.parasolCells_beta = parasolCells_beta; + _retinaParameters.IplMagno.parasolCells_tau = parasolCells_tau; + _retinaParameters.IplMagno.parasolCells_k = parasolCells_k; + _retinaParameters.IplMagno.amacrinCellsTemporalCutFrequency = amacrinCellsTemporalCutFrequency; + _retinaParameters.IplMagno.V0CompressionParameter = V0CompressionParameter; + _retinaParameters.IplMagno.localAdaptintegration_tau = localAdaptintegration_tau; + _retinaParameters.IplMagno.localAdaptintegration_k = localAdaptintegration_k; +} + +void RetinaOCLImpl::run(const InputArray input) +{ + oclMat &inputMatToConvert = getOclMatRef(input); + bool colorMode = convertToColorPlanes(inputMatToConvert, _inputBuffer); + // first convert input image to the compatible format : std::valarray + // process the retina + if (!_retinaFilter->runFilter(_inputBuffer, colorMode, false, _retinaParameters.OPLandIplParvo.colorMode && colorMode, false)) + { + throw cv::Exception(-1, "Retina cannot be applied, wrong input buffer size", "RetinaOCLImpl::run", "Retina.h", 0); + } +} + +void RetinaOCLImpl::getParvo(OutputArray output) +{ + oclMat &retinaOutput_parvo = getOclMatRef(output); + if (_retinaFilter->getColorMode()) + { + // reallocate output buffer (if necessary) + convertToInterleaved(_retinaFilter->getColorOutput(), true, retinaOutput_parvo); + } + else + { + // reallocate output buffer (if necessary) + convertToInterleaved(_retinaFilter->getContours(), false, retinaOutput_parvo); + } + //retinaOutput_parvo/=255.0; +} +void RetinaOCLImpl::getMagno(OutputArray output) +{ + oclMat &retinaOutput_magno = getOclMatRef(output); + // reallocate output buffer (if necessary) + convertToInterleaved(_retinaFilter->getMovingContours(), false, retinaOutput_magno); + //retinaOutput_magno/=255.0; +} +// private method called by constructirs +void RetinaOCLImpl::_init(const cv::Size inputSz, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght) +{ + // basic error check + if (inputSz.height*inputSz.width <= 0) + { + throw cv::Exception(-1, "Bad retina size setup : size height and with must be superior to zero", "RetinaOCLImpl::setup", "Retina.h", 0); + } + + // allocate the retina model + if (_retinaFilter) + { + delete _retinaFilter; + } + _retinaFilter = new RetinaFilter(inputSz.height, inputSz.width, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght); + + // prepare the default parameter XML file with default setup + setup(_retinaParameters); + + // init retina + _retinaFilter->clearAllBuffers(); +} + +bool RetinaOCLImpl::convertToColorPlanes(const oclMat& input, oclMat &output) +{ + oclMat convert_input; + input.convertTo(convert_input, CV_32F); + if(convert_input.channels() == 3 || convert_input.channels() == 4) + { + ocl::ensureSizeIsEnough(int(_retinaFilter->getInputNBrows() * 4), + int(_retinaFilter->getInputNBcolumns()), CV_32FC1, output); + oclMat channel_splits[4] = + { + output(Rect(Point(0, _retinaFilter->getInputNBrows() * 2), getInputSize())), + output(Rect(Point(0, _retinaFilter->getInputNBrows()), getInputSize())), + output(Rect(Point(0, 0), getInputSize())), + output(Rect(Point(0, _retinaFilter->getInputNBrows() * 3), getInputSize())) + }; + ocl::split(convert_input, channel_splits); + return true; + } + else if(convert_input.channels() == 1) + { + convert_input.copyTo(output); + return false; + } + else + { + CV_Error(-1, "Retina ocl only support 1, 3, 4 channel input"); + return false; + } +} +void RetinaOCLImpl::convertToInterleaved(const oclMat& input, bool colorMode, oclMat &output) +{ + input.convertTo(output, CV_8U); + if(colorMode) + { + int numOfSplits = input.rows / getInputSize().height; + std::vector channel_splits(numOfSplits); + for(int i = 0; i < static_cast(channel_splits.size()); i ++) + { + channel_splits[i] = + output(Rect(Point(0, _retinaFilter->getInputNBrows() * (numOfSplits - i - 1)), getInputSize())); + } + merge(channel_splits, output); + } + else + { + //... + } +} + +void RetinaOCLImpl::clearBuffers() +{ + _retinaFilter->clearAllBuffers(); +} + +void RetinaOCLImpl::activateMovingContoursProcessing(const bool activate) +{ + _retinaFilter->activateMovingContoursProcessing(activate); +} + +void RetinaOCLImpl::activateContoursProcessing(const bool activate) +{ + _retinaFilter->activateContoursProcessing(activate); +} + +/////////////////////////////////////// +///////// BasicRetinaFilter /////////// +/////////////////////////////////////// +BasicRetinaFilter::BasicRetinaFilter(const unsigned int NBrows, const unsigned int NBcolumns, const unsigned int parametersListSize, const bool) + : _NBrows(NBrows), _NBcols(NBcolumns), + _filterOutput(NBrows, NBcolumns, CV_32FC1), + _localBuffer(NBrows, NBcolumns, CV_32FC1), + _filteringCoeficientsTable(3 * parametersListSize) +{ + _halfNBrows = _filterOutput.rows / 2; + _halfNBcolumns = _filterOutput.cols / 2; + + // set default values + _maxInputValue = 256.0; + + // reset all buffers + clearAllBuffers(); +} + +BasicRetinaFilter::~BasicRetinaFilter() +{ +} + +void BasicRetinaFilter::resize(const unsigned int NBrows, const unsigned int NBcolumns) +{ + // resizing buffers + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _filterOutput); + + // updating variables + _halfNBrows = _filterOutput.rows / 2; + _halfNBcolumns = _filterOutput.cols / 2; + + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _localBuffer); + // reset buffers + clearAllBuffers(); +} + +void BasicRetinaFilter::setLPfilterParameters(const float beta, const float tau, const float desired_k, const unsigned int filterIndex) +{ + float _beta = beta + tau; + float k = desired_k; + // check if the spatial constant is correct (avoid 0 value to avoid division by 0) + if (desired_k <= 0) + { + k = 0.001f; + std::cerr << "BasicRetinaFilter::spatial constant of the low pass filter must be superior to zero !!! correcting parameter setting to 0,001" << std::endl; + } + + float _alpha = k * k; + float _mu = 0.8f; + unsigned int tableOffset = filterIndex * 3; + if (k <= 0) + { + std::cerr << "BasicRetinaFilter::spatial filtering coefficient must be superior to zero, correcting value to 0.01" << std::endl; + _alpha = 0.0001f; + } + + float _temp = (1.0f + _beta) / (2.0f * _mu * _alpha); + float a = _filteringCoeficientsTable[tableOffset] = 1.0f + _temp - (float)sqrt( (1.0f + _temp) * (1.0f + _temp) - 1.0f); + _filteringCoeficientsTable[1 + tableOffset] = (1.0f - a) * (1.0f - a) * (1.0f - a) * (1.0f - a) / (1.0f + _beta); + _filteringCoeficientsTable[2 + tableOffset] = tau; +} +const oclMat &BasicRetinaFilter::runFilter_LocalAdapdation(const oclMat &inputFrame, const oclMat &localLuminance) +{ + _localLuminanceAdaptation(inputFrame, localLuminance, _filterOutput); + return _filterOutput; +} + + +void BasicRetinaFilter::runFilter_LocalAdapdation(const oclMat &inputFrame, const oclMat &localLuminance, oclMat &outputFrame) +{ + _localLuminanceAdaptation(inputFrame, localLuminance, outputFrame); +} + +const oclMat &BasicRetinaFilter::runFilter_LocalAdapdation_autonomous(const oclMat &inputFrame) +{ + _spatiotemporalLPfilter(inputFrame, _filterOutput); + _localLuminanceAdaptation(inputFrame, _filterOutput, _filterOutput); + return _filterOutput; +} +void BasicRetinaFilter::runFilter_LocalAdapdation_autonomous(const oclMat &inputFrame, oclMat &outputFrame) +{ + _spatiotemporalLPfilter(inputFrame, _filterOutput); + _localLuminanceAdaptation(inputFrame, _filterOutput, outputFrame); +} + +void BasicRetinaFilter::_localLuminanceAdaptation(oclMat &inputOutputFrame, const oclMat &localLuminance) +{ + _localLuminanceAdaptation(inputOutputFrame, localLuminance, inputOutputFrame, false); +} + +void BasicRetinaFilter::_localLuminanceAdaptation(const oclMat &inputFrame, const oclMat &localLuminance, oclMat &outputFrame, const bool updateLuminanceMean) +{ + if (updateLuminanceMean) + { + float meanLuminance = saturate_cast(ocl::sum(inputFrame)[0]) / getNBpixels(); + updateCompressionParameter(meanLuminance); + } + int elements_per_row = static_cast(inputFrame.step / inputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, _NBrows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &localLuminance.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &inputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &_localLuminanceAddon)); + args.push_back(std::make_pair(sizeof(cl_float), &_localLuminanceFactor)); + args.push_back(std::make_pair(sizeof(cl_float), &_maxInputValue)); + openCLExecuteKernel(ctx, &retina_kernel, "localLuminanceAdaptation", globalSize, localSize, args, -1, -1); +} + +const oclMat &BasicRetinaFilter::runFilter_LPfilter(const oclMat &inputFrame, const unsigned int filterIndex) +{ + _spatiotemporalLPfilter(inputFrame, _filterOutput, filterIndex); + return _filterOutput; +} +void BasicRetinaFilter::runFilter_LPfilter(const oclMat &inputFrame, oclMat &outputFrame, const unsigned int filterIndex) +{ + _spatiotemporalLPfilter(inputFrame, outputFrame, filterIndex); +} + +void BasicRetinaFilter::_spatiotemporalLPfilter(const oclMat &inputFrame, oclMat &LPfilterOutput, const unsigned int filterIndex) +{ + unsigned int coefTableOffset = filterIndex * 3; + + _a = _filteringCoeficientsTable[coefTableOffset]; + _gain = _filteringCoeficientsTable[1 + coefTableOffset]; + _tau = _filteringCoeficientsTable[2 + coefTableOffset]; + + _horizontalCausalFilter_addInput(inputFrame, LPfilterOutput); + _horizontalAnticausalFilter(LPfilterOutput); + _verticalCausalFilter(LPfilterOutput); + _verticalAnticausalFilter_multGain(LPfilterOutput); +} + +void BasicRetinaFilter::_horizontalCausalFilter_addInput(const oclMat &inputFrame, oclMat &outputFrame) +{ + int elements_per_row = static_cast(inputFrame.step / inputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBrows, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &inputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &inputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_int), &inputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_float), &_tau)); + args.push_back(std::make_pair(sizeof(cl_float), &_a)); + openCLExecuteKernel(ctx, &retina_kernel, "horizontalCausalFilter_addInput", globalSize, localSize, args, -1, -1); +} + +void BasicRetinaFilter::_horizontalAnticausalFilter(oclMat &outputFrame) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBrows, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_float), &_a)); + openCLExecuteKernel(ctx, &retina_kernel, "horizontalAnticausalFilter", globalSize, localSize, args, -1, -1); +} + +void BasicRetinaFilter::_verticalCausalFilter(oclMat &outputFrame) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_float), &_a)); + openCLExecuteKernel(ctx, &retina_kernel, "verticalCausalFilter", globalSize, localSize, args, -1, -1); +} + +void BasicRetinaFilter::_verticalAnticausalFilter_multGain(oclMat &outputFrame) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_float), &_a)); + args.push_back(std::make_pair(sizeof(cl_float), &_gain)); + openCLExecuteKernel(ctx, &retina_kernel, "verticalAnticausalFilter_multGain", globalSize, localSize, args, -1, -1); +} + +void BasicRetinaFilter::_horizontalAnticausalFilter_Irregular(oclMat &outputFrame, const oclMat &spatialConstantBuffer) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {outputFrame.rows, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &spatialConstantBuffer.data)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_int), &spatialConstantBuffer.offset)); + openCLExecuteKernel(ctx, &retina_kernel, "horizontalAnticausalFilter_Irregular", globalSize, localSize, args, -1, -1); +} + +// vertical anticausal filter +void BasicRetinaFilter::_verticalCausalFilter_Irregular(oclMat &outputFrame, const oclMat &spatialConstantBuffer) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {outputFrame.cols, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &spatialConstantBuffer.data)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_int), &spatialConstantBuffer.offset)); + openCLExecuteKernel(ctx, &retina_kernel, "verticalCausalFilter_Irregular", globalSize, localSize, args, -1, -1); +} + +void cv::bioinspired::ocl::normalizeGrayOutput_0_maxOutputValue(oclMat &inputOutputBuffer, const float maxOutputValue) +{ + double min_val, max_val; + ocl::minMax(inputOutputBuffer, &min_val, &max_val); + float factor = maxOutputValue / static_cast(max_val - min_val); + float offset = - static_cast(min_val) * factor; + ocl::multiply(factor, inputOutputBuffer, inputOutputBuffer); + ocl::add(inputOutputBuffer, offset, inputOutputBuffer); +} + +void cv::bioinspired::ocl::normalizeGrayOutputCentredSigmoide(const float meanValue, const float sensitivity, oclMat &in, oclMat &out, const float maxValue) +{ + if (sensitivity == 1.0f) + { + std::cerr << "TemplateBuffer::TemplateBuffer::normalizeGrayOutputCentredSigmoide error: 2nd parameter (sensitivity) must not equal 0, copying original data..." << std::endl; + in.copyTo(out); + return; + } + + float X0 = maxValue / (sensitivity - 1.0f); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {in.cols, out.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + int elements_per_row = static_cast(out.step / out.elemSize()); + + args.push_back(std::make_pair(sizeof(cl_mem), &in.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &out.data)); + args.push_back(std::make_pair(sizeof(cl_int), &in.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &in.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &meanValue)); + args.push_back(std::make_pair(sizeof(cl_float), &X0)); + openCLExecuteKernel(ctx, &retina_kernel, "normalizeGrayOutputCentredSigmoide", globalSize, localSize, args, -1, -1); +} + +void cv::bioinspired::ocl::normalizeGrayOutputNearZeroCentreredSigmoide(oclMat &inputPicture, oclMat &outputBuffer, const float sensitivity, const float maxOutputValue) +{ + float X0cube = sensitivity * sensitivity * sensitivity; + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {inputPicture.cols, inputPicture.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + int elements_per_row = static_cast(inputPicture.step / inputPicture.elemSize()); + args.push_back(std::make_pair(sizeof(cl_mem), &inputPicture.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &outputBuffer.data)); + args.push_back(std::make_pair(sizeof(cl_int), &inputPicture.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &inputPicture.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &maxOutputValue)); + args.push_back(std::make_pair(sizeof(cl_float), &X0cube)); + openCLExecuteKernel(ctx, &retina_kernel, "normalizeGrayOutputNearZeroCentreredSigmoide", globalSize, localSize, args, -1, -1); +} + +void cv::bioinspired::ocl::centerReductImageLuminance(oclMat &inputoutput) +{ + Scalar mean, stddev; + cv::meanStdDev((Mat)inputoutput, mean, stddev); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {inputoutput.cols, inputoutput.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + float f_mean = static_cast(mean[0]); + float f_stddev = static_cast(stddev[0]); + int elements_per_row = static_cast(inputoutput.step / inputoutput.elemSize()); + args.push_back(std::make_pair(sizeof(cl_mem), &inputoutput.data)); + args.push_back(std::make_pair(sizeof(cl_int), &inputoutput.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &inputoutput.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &f_mean)); + args.push_back(std::make_pair(sizeof(cl_float), &f_stddev)); + openCLExecuteKernel(ctx, &retina_kernel, "centerReductImageLuminance", globalSize, localSize, args, -1, -1); +} + +/////////////////////////////////////// +///////// ParvoRetinaFilter /////////// +/////////////////////////////////////// +ParvoRetinaFilter::ParvoRetinaFilter(const unsigned int NBrows, const unsigned int NBcolumns) + : BasicRetinaFilter(NBrows, NBcolumns, 3), + _photoreceptorsOutput(NBrows, NBcolumns, CV_32FC1), + _horizontalCellsOutput(NBrows, NBcolumns, CV_32FC1), + _parvocellularOutputON(NBrows, NBcolumns, CV_32FC1), + _parvocellularOutputOFF(NBrows, NBcolumns, CV_32FC1), + _bipolarCellsOutputON(NBrows, NBcolumns, CV_32FC1), + _bipolarCellsOutputOFF(NBrows, NBcolumns, CV_32FC1), + _localAdaptationOFF(NBrows, NBcolumns, CV_32FC1) +{ + // link to the required local parent adaptation buffers + _localAdaptationON = _localBuffer; + _parvocellularOutputONminusOFF = _filterOutput; + + // init: set all the values to 0 + clearAllBuffers(); +} + +ParvoRetinaFilter::~ParvoRetinaFilter() +{ +} + +void ParvoRetinaFilter::clearAllBuffers() +{ + BasicRetinaFilter::clearAllBuffers(); + _photoreceptorsOutput = 0; + _horizontalCellsOutput = 0; + _parvocellularOutputON = 0; + _parvocellularOutputOFF = 0; + _bipolarCellsOutputON = 0; + _bipolarCellsOutputOFF = 0; + _localAdaptationOFF = 0; +} +void ParvoRetinaFilter::resize(const unsigned int NBrows, const unsigned int NBcolumns) +{ + BasicRetinaFilter::resize(NBrows, NBcolumns); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _photoreceptorsOutput); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _horizontalCellsOutput); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _parvocellularOutputON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _parvocellularOutputOFF); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _bipolarCellsOutputON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _bipolarCellsOutputOFF); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _localAdaptationOFF); + + // link to the required local parent adaptation buffers + _localAdaptationON = _localBuffer; + _parvocellularOutputONminusOFF = _filterOutput; + + // clean buffers + clearAllBuffers(); +} + +void ParvoRetinaFilter::setOPLandParvoFiltersParameters(const float beta1, const float tau1, const float k1, const float beta2, const float tau2, const float k2) +{ + // init photoreceptors low pass filter + setLPfilterParameters(beta1, tau1, k1); + // init horizontal cells low pass filter + setLPfilterParameters(beta2, tau2, k2, 1); + // init parasol ganglion cells low pass filter (default parameters) + setLPfilterParameters(0, tau1, k1, 2); + +} +const oclMat &ParvoRetinaFilter::runFilter(const oclMat &inputFrame, const bool useParvoOutput) +{ + _spatiotemporalLPfilter(inputFrame, _photoreceptorsOutput); + _spatiotemporalLPfilter(_photoreceptorsOutput, _horizontalCellsOutput, 1); + _OPL_OnOffWaysComputing(); + + if (useParvoOutput) + { + // local adaptation processes on ON and OFF ways + _spatiotemporalLPfilter(_bipolarCellsOutputON, _localAdaptationON, 2); + _localLuminanceAdaptation(_parvocellularOutputON, _localAdaptationON); + _spatiotemporalLPfilter(_bipolarCellsOutputOFF, _localAdaptationOFF, 2); + _localLuminanceAdaptation(_parvocellularOutputOFF, _localAdaptationOFF); + ocl::subtract(_parvocellularOutputON, _parvocellularOutputOFF, _parvocellularOutputONminusOFF); + } + + return _parvocellularOutputONminusOFF; +} +void ParvoRetinaFilter::_OPL_OnOffWaysComputing() +{ + int elements_per_row = static_cast(_photoreceptorsOutput.step / _photoreceptorsOutput.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {(_photoreceptorsOutput.cols + 3) / 4, _photoreceptorsOutput.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &_photoreceptorsOutput.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_horizontalCellsOutput.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_bipolarCellsOutputON.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_bipolarCellsOutputOFF.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_parvocellularOutputON.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_parvocellularOutputOFF.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_photoreceptorsOutput.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &_photoreceptorsOutput.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "OPL_OnOffWaysComputing", globalSize, localSize, args, -1, -1); +} + +/////////////////////////////////////// +//////////// MagnoFilter ////////////// +/////////////////////////////////////// +MagnoRetinaFilter::MagnoRetinaFilter(const unsigned int NBrows, const unsigned int NBcolumns) + : BasicRetinaFilter(NBrows, NBcolumns, 2), + _previousInput_ON(NBrows, NBcolumns, CV_32FC1), + _previousInput_OFF(NBrows, NBcolumns, CV_32FC1), + _amacrinCellsTempOutput_ON(NBrows, NBcolumns, CV_32FC1), + _amacrinCellsTempOutput_OFF(NBrows, NBcolumns, CV_32FC1), + _magnoXOutputON(NBrows, NBcolumns, CV_32FC1), + _magnoXOutputOFF(NBrows, NBcolumns, CV_32FC1), + _localProcessBufferON(NBrows, NBcolumns, CV_32FC1), + _localProcessBufferOFF(NBrows, NBcolumns, CV_32FC1) +{ + _magnoYOutput = _filterOutput; + _magnoYsaturated = _localBuffer; + + clearAllBuffers(); +} + +MagnoRetinaFilter::~MagnoRetinaFilter() +{ +} +void MagnoRetinaFilter::clearAllBuffers() +{ + BasicRetinaFilter::clearAllBuffers(); + _previousInput_ON = 0; + _previousInput_OFF = 0; + _amacrinCellsTempOutput_ON = 0; + _amacrinCellsTempOutput_OFF = 0; + _magnoXOutputON = 0; + _magnoXOutputOFF = 0; + _localProcessBufferON = 0; + _localProcessBufferOFF = 0; + +} +void MagnoRetinaFilter::resize(const unsigned int NBrows, const unsigned int NBcolumns) +{ + BasicRetinaFilter::resize(NBrows, NBcolumns); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _previousInput_ON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _previousInput_OFF); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _amacrinCellsTempOutput_ON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _amacrinCellsTempOutput_OFF); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _magnoXOutputON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _magnoXOutputOFF); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _localProcessBufferON); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _localProcessBufferOFF); + + // to be sure, relink buffers + _magnoYOutput = _filterOutput; + _magnoYsaturated = _localBuffer; + + // reset all buffers + clearAllBuffers(); +} + +void MagnoRetinaFilter::setCoefficientsTable(const float parasolCells_beta, const float parasolCells_tau, const float parasolCells_k, const float amacrinCellsTemporalCutFrequency, const float localAdaptIntegration_tau, const float localAdaptIntegration_k ) +{ + _temporalCoefficient = (float)std::exp(-1.0f / amacrinCellsTemporalCutFrequency); + // the first set of parameters is dedicated to the low pass filtering property of the ganglion cells + BasicRetinaFilter::setLPfilterParameters(parasolCells_beta, parasolCells_tau, parasolCells_k, 0); + // the second set of parameters is dedicated to the ganglion cells output intergartion for their local adaptation property + BasicRetinaFilter::setLPfilterParameters(0, localAdaptIntegration_tau, localAdaptIntegration_k, 1); +} + +void MagnoRetinaFilter::_amacrineCellsComputing( + const oclMat &OPL_ON, + const oclMat &OPL_OFF +) +{ + int elements_per_row = static_cast(OPL_ON.step / OPL_ON.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {OPL_ON.cols, OPL_ON.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &OPL_ON.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &OPL_OFF.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_previousInput_ON.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_previousInput_OFF.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_amacrinCellsTempOutput_ON.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &_amacrinCellsTempOutput_OFF.data)); + args.push_back(std::make_pair(sizeof(cl_int), &OPL_ON.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &OPL_ON.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &_temporalCoefficient)); + openCLExecuteKernel(ctx, &retina_kernel, "amacrineCellsComputing", globalSize, localSize, args, -1, -1); +} + +const oclMat &MagnoRetinaFilter::runFilter(const oclMat &OPL_ON, const oclMat &OPL_OFF) +{ + // Compute the high pass temporal filter + _amacrineCellsComputing(OPL_ON, OPL_OFF); + + // apply low pass filtering on ON and OFF ways after temporal high pass filtering + _spatiotemporalLPfilter(_amacrinCellsTempOutput_ON, _magnoXOutputON, 0); + _spatiotemporalLPfilter(_amacrinCellsTempOutput_OFF, _magnoXOutputOFF, 0); + + // local adaptation of the ganglion cells to the local contrast of the moving contours + _spatiotemporalLPfilter(_magnoXOutputON, _localProcessBufferON, 1); + _localLuminanceAdaptation(_magnoXOutputON, _localProcessBufferON); + + _spatiotemporalLPfilter(_magnoXOutputOFF, _localProcessBufferOFF, 1); + _localLuminanceAdaptation(_magnoXOutputOFF, _localProcessBufferOFF); + + _magnoYOutput = _magnoXOutputON + _magnoXOutputOFF; + + return _magnoYOutput; +} + +/////////////////////////////////////// +//////////// RetinaColor ////////////// +/////////////////////////////////////// + +// define an array of ROI headers of input x +#define MAKE_OCLMAT_SLICES(x, n) \ + oclMat x##_slices[n];\ + for(int _SLICE_INDEX_ = 0; _SLICE_INDEX_ < n; _SLICE_INDEX_ ++)\ + {\ + x##_slices[_SLICE_INDEX_] = x(getROI(_SLICE_INDEX_));\ + } + +static float _LMStoACr1Cr2[] = {1.0, 1.0, 0.0, 1.0, -1.0, 0.0, -0.5, -0.5, 1.0}; +static float _LMStoLab[] = {0.5774f, 0.5774f, 0.5774f, 0.4082f, 0.4082f, -0.8165f, 0.7071f, -0.7071f, 0.f}; + +RetinaColor::RetinaColor(const unsigned int NBrows, const unsigned int NBcolumns, const int samplingMethod) + : BasicRetinaFilter(NBrows, NBcolumns, 3), + _RGBmosaic(NBrows * 3, NBcolumns, CV_32FC1), + _tempMultiplexedFrame(NBrows, NBcolumns, CV_32FC1), + _demultiplexedTempBuffer(NBrows * 3, NBcolumns, CV_32FC1), + _demultiplexedColorFrame(NBrows * 3, NBcolumns, CV_32FC1), + _chrominance(NBrows * 3, NBcolumns, CV_32FC1), + _colorLocalDensity(NBrows * 3, NBcolumns, CV_32FC1), + _imageGradient(NBrows * 3, NBcolumns, CV_32FC1) +{ + // link to parent buffers (let's recycle !) + _luminance = _filterOutput; + _multiplexedFrame = _localBuffer; + + _objectInit = false; + _samplingMethod = samplingMethod; + _saturateColors = false; + _colorSaturationValue = 4.0; + + // set default spatio-temporal filter parameters + setLPfilterParameters(0.0, 0.0, 1.5); + setLPfilterParameters(0.0, 0.0, 10.5, 1);// for the low pass filter dedicated to contours energy extraction (demultiplexing process) + setLPfilterParameters(0.f, 0.f, 0.9f, 2); + + // init default value on image Gradient + _imageGradient = 0.57f; + + // init color sampling map + _initColorSampling(); + + // flush all buffers + clearAllBuffers(); +} + +RetinaColor::~RetinaColor() +{ + +} + +void RetinaColor::clearAllBuffers() +{ + BasicRetinaFilter::clearAllBuffers(); + _tempMultiplexedFrame = 0.f; + _demultiplexedTempBuffer = 0.f; + + _demultiplexedColorFrame = 0.f; + _chrominance = 0.f; + _imageGradient = 0.57f; +} + +void RetinaColor::resize(const unsigned int NBrows, const unsigned int NBcolumns) +{ + BasicRetinaFilter::clearAllBuffers(); + ensureSizeIsEnough(NBrows, NBcolumns, CV_32FC1, _tempMultiplexedFrame); + ensureSizeIsEnough(NBrows * 2, NBcolumns, CV_32FC1, _imageGradient); + ensureSizeIsEnough(NBrows * 3, NBcolumns, CV_32FC1, _RGBmosaic); + ensureSizeIsEnough(NBrows * 3, NBcolumns, CV_32FC1, _demultiplexedTempBuffer); + ensureSizeIsEnough(NBrows * 3, NBcolumns, CV_32FC1, _demultiplexedColorFrame); + ensureSizeIsEnough(NBrows * 3, NBcolumns, CV_32FC1, _chrominance); + ensureSizeIsEnough(NBrows * 3, NBcolumns, CV_32FC1, _colorLocalDensity); + + // link to parent buffers (let's recycle !) + _luminance = _filterOutput; + _multiplexedFrame = _localBuffer; + + // init color sampling map + _initColorSampling(); + + // clean buffers + clearAllBuffers(); +} + +static void inverseValue(oclMat &input) +{ + int elements_per_row = static_cast(input.step / input.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {input.cols, input.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &input.data)); + args.push_back(std::make_pair(sizeof(cl_int), &input.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &input.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "inverseValue", globalSize, localSize, args, -1, -1); +} + +void RetinaColor::_initColorSampling() +{ + CV_Assert(_samplingMethod == RETINA_COLOR_BAYER); + _pR = _pB = 0.25; + _pG = 0.5; + // filling the mosaic buffer: + _RGBmosaic = 0; + Mat tmp_mat(_NBrows * 3, _NBcols, CV_32FC1); + float * tmp_mat_ptr = tmp_mat.ptr(); + tmp_mat.setTo(0); + for (unsigned int index = 0 ; index < getNBpixels(); ++index) + { + tmp_mat_ptr[bayerSampleOffset(index)] = 1.0; + } + _RGBmosaic.upload(tmp_mat); + // computing photoreceptors local density + MAKE_OCLMAT_SLICES(_RGBmosaic, 3); + MAKE_OCLMAT_SLICES(_colorLocalDensity, 3); + + _spatiotemporalLPfilter(_RGBmosaic_slices[0], _colorLocalDensity_slices[0]); + _spatiotemporalLPfilter(_RGBmosaic_slices[1], _colorLocalDensity_slices[1]); + _spatiotemporalLPfilter(_RGBmosaic_slices[2], _colorLocalDensity_slices[2]); + + //_colorLocalDensity = oclMat(_colorLocalDensity.size(), _colorLocalDensity.type(), 1.f) / _colorLocalDensity; + inverseValue(_colorLocalDensity); + + _objectInit = true; +} + +static void demultiplex(const oclMat &input, oclMat &ouput) +{ + int elements_per_row = static_cast(input.step / input.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {input.cols, input.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &input.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &ouput.data)); + args.push_back(std::make_pair(sizeof(cl_int), &input.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &input.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "runColorDemultiplexingBayer", globalSize, localSize, args, -1, -1); +} + +static void normalizePhotoDensity( + const oclMat &chroma, + const oclMat &colorDensity, + const oclMat &multiplex, + oclMat &ocl_luma, + oclMat &demultiplex, + const float pG +) +{ + int elements_per_row = static_cast(ocl_luma.step / ocl_luma.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {ocl_luma.cols, ocl_luma.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &chroma.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &colorDensity.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &multiplex.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &ocl_luma.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &demultiplex.data)); + args.push_back(std::make_pair(sizeof(cl_int), &ocl_luma.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &ocl_luma.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &pG)); + openCLExecuteKernel(ctx, &retina_kernel, "normalizePhotoDensity", globalSize, localSize, args, -1, -1); +} + +static void substractResidual( + oclMat &colorDemultiplex, + float pR, + float pG, + float pB +) +{ + int elements_per_row = static_cast(colorDemultiplex.step / colorDemultiplex.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + int rows = colorDemultiplex.rows / 3, cols = colorDemultiplex.cols; + size_t globalSize[] = {cols, rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &colorDemultiplex.data)); + args.push_back(std::make_pair(sizeof(cl_int), &cols)); + args.push_back(std::make_pair(sizeof(cl_int), &rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &pR)); + args.push_back(std::make_pair(sizeof(cl_float), &pG)); + args.push_back(std::make_pair(sizeof(cl_float), &pB)); + openCLExecuteKernel(ctx, &retina_kernel, "substractResidual", globalSize, localSize, args, -1, -1); +} + +static void demultiplexAssign(const oclMat& input, const oclMat& output) +{ + // only supports bayer + int elements_per_row = static_cast(input.step / input.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + int rows = input.rows / 3, cols = input.cols; + size_t globalSize[] = {cols, rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &input.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &output.data)); + args.push_back(std::make_pair(sizeof(cl_int), &cols)); + args.push_back(std::make_pair(sizeof(cl_int), &rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "demultiplexAssign", globalSize, localSize, args, -1, -1); +} + +void RetinaColor::runColorDemultiplexing( + const oclMat &ocl_multiplexed_input, + const bool adaptiveFiltering, + const float maxInputValue +) +{ + MAKE_OCLMAT_SLICES(_demultiplexedTempBuffer, 3); + MAKE_OCLMAT_SLICES(_chrominance, 3); + MAKE_OCLMAT_SLICES(_RGBmosaic, 3); + MAKE_OCLMAT_SLICES(_demultiplexedColorFrame, 3); + MAKE_OCLMAT_SLICES(_colorLocalDensity, 3); + + _demultiplexedTempBuffer.setTo(0); + demultiplex(ocl_multiplexed_input, _demultiplexedTempBuffer); + + // interpolate the demultiplexed frame depending on the color sampling method + if (!adaptiveFiltering) + { + CV_Assert(adaptiveFiltering == false); + } + + _spatiotemporalLPfilter(_demultiplexedTempBuffer_slices[0], _chrominance_slices[0]); + _spatiotemporalLPfilter(_demultiplexedTempBuffer_slices[1], _chrominance_slices[1]); + _spatiotemporalLPfilter(_demultiplexedTempBuffer_slices[2], _chrominance_slices[2]); + + if (!adaptiveFiltering)// compute the gradient on the luminance + { + // TODO: implement me! + CV_Assert(adaptiveFiltering == false); + } + else + { + normalizePhotoDensity(_chrominance, _colorLocalDensity, ocl_multiplexed_input, _luminance, _demultiplexedTempBuffer, _pG); + // compute the gradient of the luminance + _computeGradient(_luminance, _imageGradient); + + _adaptiveSpatialLPfilter(_RGBmosaic_slices[0], _imageGradient, _chrominance_slices[0]); + _adaptiveSpatialLPfilter(_RGBmosaic_slices[1], _imageGradient, _chrominance_slices[1]); + _adaptiveSpatialLPfilter(_RGBmosaic_slices[2], _imageGradient, _chrominance_slices[2]); + + _adaptiveSpatialLPfilter(_demultiplexedTempBuffer_slices[0], _imageGradient, _demultiplexedColorFrame_slices[0]); + _adaptiveSpatialLPfilter(_demultiplexedTempBuffer_slices[1], _imageGradient, _demultiplexedColorFrame_slices[1]); + _adaptiveSpatialLPfilter(_demultiplexedTempBuffer_slices[2], _imageGradient, _demultiplexedColorFrame_slices[2]); + + _demultiplexedColorFrame /= _chrominance; // per element division + substractResidual(_demultiplexedColorFrame, _pR, _pG, _pB); + runColorMultiplexing(_demultiplexedColorFrame, _tempMultiplexedFrame); + + _demultiplexedTempBuffer.setTo(0); + _luminance = ocl_multiplexed_input - _tempMultiplexedFrame; + demultiplexAssign(_demultiplexedColorFrame, _demultiplexedTempBuffer); + + for(int i = 0; i < 3; i ++) + { + _spatiotemporalLPfilter(_demultiplexedTempBuffer_slices[i], _demultiplexedTempBuffer_slices[i]); + _demultiplexedColorFrame_slices[i] = _demultiplexedTempBuffer_slices[i] * _colorLocalDensity_slices[i] + _luminance; + } + } + // eliminate saturated colors by simple clipping values to the input range + clipRGBOutput_0_maxInputValue(_demultiplexedColorFrame, maxInputValue); + + if (_saturateColors) + { + ocl::normalizeGrayOutputCentredSigmoide(128, maxInputValue, _demultiplexedColorFrame, _demultiplexedColorFrame); + } +} +void RetinaColor::runColorMultiplexing(const oclMat &demultiplexedInputFrame, oclMat &multiplexedFrame) +{ + int elements_per_row = static_cast(multiplexedFrame.step / multiplexedFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {multiplexedFrame.cols, multiplexedFrame.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &demultiplexedInputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &multiplexedFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &multiplexedFrame.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &multiplexedFrame.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "runColorMultiplexingBayer", globalSize, localSize, args, -1, -1); +} + +void RetinaColor::clipRGBOutput_0_maxInputValue(oclMat &inputOutputBuffer, const float maxInputValue) +{ + // the kernel is equivalent to: + //ocl::threshold(inputOutputBuffer, inputOutputBuffer, maxInputValue, maxInputValue, THRESH_TRUNC); + //ocl::threshold(inputOutputBuffer, inputOutputBuffer, 0, 0, THRESH_TOZERO); + int elements_per_row = static_cast(inputOutputBuffer.step / inputOutputBuffer.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, inputOutputBuffer.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &inputOutputBuffer.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &inputOutputBuffer.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &maxInputValue)); + openCLExecuteKernel(ctx, &retina_kernel, "clipRGBOutput_0_maxInputValue", globalSize, localSize, args, -1, -1); +} + +void RetinaColor::_adaptiveSpatialLPfilter(const oclMat &inputFrame, const oclMat &gradient, oclMat &outputFrame) +{ + /**********/ + _gain = (1 - 0.57f) * (1 - 0.57f) * (1 - 0.06f) * (1 - 0.06f); + + // launch the serie of 1D directional filters in order to compute the 2D low pass filter + // -> horizontal filters work with the first layer of imageGradient + _adaptiveHorizontalCausalFilter_addInput(inputFrame, gradient, outputFrame); + _horizontalAnticausalFilter_Irregular(outputFrame, gradient); + // -> horizontal filters work with the second layer of imageGradient + _verticalCausalFilter_Irregular(outputFrame, gradient(getROI(1))); + _adaptiveVerticalAnticausalFilter_multGain(gradient, outputFrame); +} + +void RetinaColor::_adaptiveHorizontalCausalFilter_addInput(const oclMat &inputFrame, const oclMat &gradient, oclMat &outputFrame) +{ + int elements_per_row = static_cast(inputFrame.step / inputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBrows, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &inputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &gradient.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &inputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_int), &gradient.offset)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + openCLExecuteKernel(ctx, &retina_kernel, "adaptiveHorizontalCausalFilter_addInput", globalSize, localSize, args, -1, -1); +} + +void RetinaColor::_adaptiveVerticalAnticausalFilter_multGain(const oclMat &gradient, oclMat &outputFrame) +{ + int elements_per_row = static_cast(outputFrame.step / outputFrame.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, 1, 1}; + size_t localSize[] = {256, 1, 1}; + + int gradOffset = gradient.offset + static_cast(gradient.step * _NBrows); + + args.push_back(std::make_pair(sizeof(cl_mem), &gradient.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &outputFrame.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_int), &gradOffset)); + args.push_back(std::make_pair(sizeof(cl_int), &outputFrame.offset)); + args.push_back(std::make_pair(sizeof(cl_float), &_gain)); + openCLExecuteKernel(ctx, &retina_kernel, "adaptiveVerticalAnticausalFilter_multGain", globalSize, localSize, args, -1, -1); +} +void RetinaColor::_computeGradient(const oclMat &luminance, oclMat &gradient) +{ + int elements_per_row = static_cast(luminance.step / luminance.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {_NBcols, _NBrows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &luminance.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &gradient.data)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBcols)); + args.push_back(std::make_pair(sizeof(cl_int), &_NBrows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + openCLExecuteKernel(ctx, &retina_kernel, "computeGradient", globalSize, localSize, args, -1, -1); +} + +/////////////////////////////////////// +//////////// RetinaFilter ///////////// +/////////////////////////////////////// +RetinaFilter::RetinaFilter(const unsigned int sizeRows, const unsigned int sizeColumns, const bool colorMode, const int samplingMethod, const bool useRetinaLogSampling, const double, const double) + : + _photoreceptorsPrefilter(sizeRows, sizeColumns, 4), + _ParvoRetinaFilter(sizeRows, sizeColumns), + _MagnoRetinaFilter(sizeRows, sizeColumns), + _colorEngine(sizeRows, sizeColumns, samplingMethod) +{ + CV_Assert(!useRetinaLogSampling); + + // set default processing activities + _useParvoOutput = true; + _useMagnoOutput = true; + + _useColorMode = colorMode; + + // set default parameters + setGlobalParameters(); + + // stability controls values init + _setInitPeriodCount(); + _globalTemporalConstant = 25; + + // reset all buffers + clearAllBuffers(); +} + +RetinaFilter::~RetinaFilter() +{ +} + +void RetinaFilter::clearAllBuffers() +{ + _photoreceptorsPrefilter.clearAllBuffers(); + _ParvoRetinaFilter.clearAllBuffers(); + _MagnoRetinaFilter.clearAllBuffers(); + _colorEngine.clearAllBuffers(); + // stability controls value init + _setInitPeriodCount(); +} + +void RetinaFilter::resize(const unsigned int NBrows, const unsigned int NBcolumns) +{ + unsigned int rows = NBrows, cols = NBcolumns; + + // resize optionnal member and adjust other modules size if required + _photoreceptorsPrefilter.resize(rows, cols); + _ParvoRetinaFilter.resize(rows, cols); + _MagnoRetinaFilter.resize(rows, cols); + _colorEngine.resize(rows, cols); + + // clean buffers + clearAllBuffers(); + +} + +void RetinaFilter::_setInitPeriodCount() +{ + // find out the maximum temporal constant value and apply a security factor + // false value (obviously too long) but appropriate for simple use + _globalTemporalConstant = (unsigned int)(_ParvoRetinaFilter.getPhotoreceptorsTemporalConstant() + _ParvoRetinaFilter.getHcellsTemporalConstant() + _MagnoRetinaFilter.getTemporalConstant()); + // reset frame counter + _ellapsedFramesSinceLastReset = 0; +} + +void RetinaFilter::setGlobalParameters(const float OPLspatialResponse1, const float OPLtemporalresponse1, const float OPLassymetryGain, const float OPLspatialResponse2, const float OPLtemporalresponse2, const float LPfilterSpatialResponse, const float LPfilterGain, const float LPfilterTemporalresponse, const float MovingContoursExtractorCoefficient, const bool normalizeParvoOutput_0_maxOutputValue, const bool normalizeMagnoOutput_0_maxOutputValue, const float maxOutputValue, const float maxInputValue, const float meanValue) +{ + _normalizeParvoOutput_0_maxOutputValue = normalizeParvoOutput_0_maxOutputValue; + _normalizeMagnoOutput_0_maxOutputValue = normalizeMagnoOutput_0_maxOutputValue; + _maxOutputValue = maxOutputValue; + _photoreceptorsPrefilter.setV0CompressionParameter(0.9f, maxInputValue, meanValue); + _photoreceptorsPrefilter.setLPfilterParameters(0, 0, 10, 3); // keeps low pass filter with low cut frequency in memory (usefull for the tone mapping function) + _ParvoRetinaFilter.setOPLandParvoFiltersParameters(0, OPLtemporalresponse1, OPLspatialResponse1, OPLassymetryGain, OPLtemporalresponse2, OPLspatialResponse2); + _ParvoRetinaFilter.setV0CompressionParameter(0.9f, maxInputValue, meanValue); + _MagnoRetinaFilter.setCoefficientsTable(LPfilterGain, LPfilterTemporalresponse, LPfilterSpatialResponse, MovingContoursExtractorCoefficient, 0, 2.0f * LPfilterSpatialResponse); + _MagnoRetinaFilter.setV0CompressionParameter(0.7f, maxInputValue, meanValue); + + // stability controls value init + _setInitPeriodCount(); +} + +bool RetinaFilter::checkInput(const oclMat &input, const bool) +{ + BasicRetinaFilter *inputTarget = &_photoreceptorsPrefilter; + + bool test = (input.rows == static_cast(inputTarget->getNBrows()) + || input.rows == static_cast(inputTarget->getNBrows()) * 3 + || input.rows == static_cast(inputTarget->getNBrows()) * 4) + && input.cols == static_cast(inputTarget->getNBcolumns()); + if (!test) + { + std::cerr << "RetinaFilter::checkInput: input buffer does not match retina buffer size, conversion aborted" << std::endl; + return false; + } + + return true; +} + +// main function that runs the filter for a given input frame +bool RetinaFilter::runFilter(const oclMat &imageInput, const bool useAdaptiveFiltering, const bool processRetinaParvoMagnoMapping, const bool useColorMode, const bool inputIsColorMultiplexed) +{ + // preliminary check + bool processSuccess = true; + if (!checkInput(imageInput, useColorMode)) + { + return false; + } + + // run the color multiplexing if needed and compute each suub filter of the retina: + // -> local adaptation + // -> contours OPL extraction + // -> moving contours extraction + + // stability controls value update + ++_ellapsedFramesSinceLastReset; + + _useColorMode = useColorMode; + + oclMat selectedPhotoreceptorsLocalAdaptationInput = imageInput; + oclMat selectedPhotoreceptorsColorInput = imageInput; + + //********** Following is input data specific photoreceptors processing + if (useColorMode && (!inputIsColorMultiplexed)) // not multiplexed color input case + { + _colorEngine.runColorMultiplexing(selectedPhotoreceptorsColorInput); + selectedPhotoreceptorsLocalAdaptationInput = _colorEngine.getMultiplexedFrame(); + } + //********** Following is generic Retina processing + + // photoreceptors local adaptation + _photoreceptorsPrefilter.runFilter_LocalAdapdation(selectedPhotoreceptorsLocalAdaptationInput, _ParvoRetinaFilter.getHorizontalCellsOutput()); + + // run parvo filter + _ParvoRetinaFilter.runFilter(_photoreceptorsPrefilter.getOutput(), _useParvoOutput); + + if (_useParvoOutput) + { + _ParvoRetinaFilter.normalizeGrayOutputCentredSigmoide(); // models the saturation of the cells, usefull for visualisation of the ON-OFF Parvo Output, Bipolar cells outputs do not change !!! + _ParvoRetinaFilter.centerReductImageLuminance(); // best for further spectrum analysis + + if (_normalizeParvoOutput_0_maxOutputValue) + { + _ParvoRetinaFilter.normalizeGrayOutput_0_maxOutputValue(_maxOutputValue); + } + } + + if (_useParvoOutput && _useMagnoOutput) + { + _MagnoRetinaFilter.runFilter(_ParvoRetinaFilter.getBipolarCellsON(), _ParvoRetinaFilter.getBipolarCellsOFF()); + if (_normalizeMagnoOutput_0_maxOutputValue) + { + _MagnoRetinaFilter.normalizeGrayOutput_0_maxOutputValue(_maxOutputValue); + } + _MagnoRetinaFilter.normalizeGrayOutputNearZeroCentreredSigmoide(); + } + + if (_useParvoOutput && _useMagnoOutput && processRetinaParvoMagnoMapping) + { + _processRetinaParvoMagnoMapping(); + if (_useColorMode) + { + _colorEngine.runColorDemultiplexing(_retinaParvoMagnoMappedFrame, useAdaptiveFiltering, _maxOutputValue); + } + return processSuccess; + } + + if (_useParvoOutput && _useColorMode) + { + _colorEngine.runColorDemultiplexing(_ParvoRetinaFilter.getOutput(), useAdaptiveFiltering, _maxOutputValue); + } + return processSuccess; +} + +const oclMat &RetinaFilter::getContours() +{ + if (_useColorMode) + { + return _colorEngine.getLuminance(); + } + else + { + return _ParvoRetinaFilter.getOutput(); + } +} +void RetinaFilter::_processRetinaParvoMagnoMapping() +{ + oclMat parvo = _ParvoRetinaFilter.getOutput(); + oclMat magno = _MagnoRetinaFilter.getOutput(); + + int halfRows = parvo.rows / 2; + int halfCols = parvo.cols / 2; + float minDistance = MIN(halfRows, halfCols) * 0.7f; + + int elements_per_row = static_cast(parvo.step / parvo.elemSize()); + + Context * ctx = Context::getContext(); + std::vector > args; + size_t globalSize[] = {parvo.cols, parvo.rows, 1}; + size_t localSize[] = {16, 16, 1}; + + args.push_back(std::make_pair(sizeof(cl_mem), &parvo.data)); + args.push_back(std::make_pair(sizeof(cl_mem), &magno.data)); + args.push_back(std::make_pair(sizeof(cl_int), &parvo.cols)); + args.push_back(std::make_pair(sizeof(cl_int), &parvo.rows)); + args.push_back(std::make_pair(sizeof(cl_int), &halfCols)); + args.push_back(std::make_pair(sizeof(cl_int), &halfRows)); + args.push_back(std::make_pair(sizeof(cl_int), &elements_per_row)); + args.push_back(std::make_pair(sizeof(cl_float), &minDistance)); + openCLExecuteKernel(ctx, &retina_kernel, "processRetinaParvoMagnoMapping", globalSize, localSize, args, -1, -1); +} +} /* namespace ocl */ + +Ptr createRetina_OCL(Size getInputSize){ return new ocl::RetinaOCLImpl(getInputSize); } +Ptr createRetina_OCL(Size getInputSize, const bool colorMode, int colorSamplingMethod, const bool useRetinaLogSampling, const double reductionFactor, const double samplingStrenght) +{ + return new ocl::RetinaOCLImpl(getInputSize, colorMode, colorSamplingMethod, useRetinaLogSampling, reductionFactor, samplingStrenght); +} + +} /* namespace bioinspired */ +} /* namespace cv */ + +#endif /* #ifdef HAVE_OPENCV_OCL */ diff --git a/modules/bioinspired/src/retina_ocl.hpp b/modules/bioinspired/src/retina_ocl.hpp new file mode 100644 index 0000000..69a833c --- /dev/null +++ b/modules/bioinspired/src/retina_ocl.hpp @@ -0,0 +1,633 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#ifndef __OCL_RETINA_HPP__ +#define __OCL_RETINA_HPP__ + +#include "precomp.hpp" + +#ifdef HAVE_OPENCV_OCL + +// please refer to c++ headers for API comments +namespace cv +{ +namespace bioinspired +{ +namespace ocl +{ +void normalizeGrayOutputCentredSigmoide(const float meanValue, const float sensitivity, cv::ocl::oclMat &in, cv::ocl::oclMat &out, const float maxValue = 255.f); +void normalizeGrayOutput_0_maxOutputValue(cv::ocl::oclMat &inputOutputBuffer, const float maxOutputValue = 255.0); +void normalizeGrayOutputNearZeroCentreredSigmoide(cv::ocl::oclMat &inputPicture, cv::ocl::oclMat &outputBuffer, const float sensitivity = 40, const float maxOutputValue = 255.0f); +void centerReductImageLuminance(cv::ocl::oclMat &inputOutputBuffer); + +class BasicRetinaFilter +{ +public: + BasicRetinaFilter(const unsigned int NBrows, const unsigned int NBcolumns, const unsigned int parametersListSize = 1, const bool useProgressiveFilter = false); + ~BasicRetinaFilter(); + inline void clearOutputBuffer() + { + _filterOutput = 0; + }; + inline void clearSecondaryBuffer() + { + _localBuffer = 0; + }; + inline void clearAllBuffers() + { + clearOutputBuffer(); + clearSecondaryBuffer(); + }; + void resize(const unsigned int NBrows, const unsigned int NBcolumns); + const cv::ocl::oclMat &runFilter_LPfilter(const cv::ocl::oclMat &inputFrame, const unsigned int filterIndex = 0); + void runFilter_LPfilter(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame, const unsigned int filterIndex = 0); + void runFilter_LPfilter_Autonomous(cv::ocl::oclMat &inputOutputFrame, const unsigned int filterIndex = 0); + const cv::ocl::oclMat &runFilter_LocalAdapdation(const cv::ocl::oclMat &inputOutputFrame, const cv::ocl::oclMat &localLuminance); + void runFilter_LocalAdapdation(const cv::ocl::oclMat &inputFrame, const cv::ocl::oclMat &localLuminance, cv::ocl::oclMat &outputFrame); + const cv::ocl::oclMat &runFilter_LocalAdapdation_autonomous(const cv::ocl::oclMat &inputFrame); + void runFilter_LocalAdapdation_autonomous(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame); + void setLPfilterParameters(const float beta, const float tau, const float k, const unsigned int filterIndex = 0); + inline void setV0CompressionParameter(const float v0, const float maxInputValue, const float) + { + _v0 = v0 * maxInputValue; + _localLuminanceFactor = v0; + _localLuminanceAddon = maxInputValue * (1.0f - v0); + _maxInputValue = maxInputValue; + }; + inline void setV0CompressionParameter(const float v0, const float meanLuminance) + { + this->setV0CompressionParameter(v0, _maxInputValue, meanLuminance); + }; + inline void setV0CompressionParameter(const float v0) + { + _v0 = v0 * _maxInputValue; + _localLuminanceFactor = v0; + _localLuminanceAddon = _maxInputValue * (1.0f - v0); + }; + inline void setV0CompressionParameterToneMapping(const float v0, const float maxInputValue, const float meanLuminance = 128.0f) + { + _v0 = v0 * maxInputValue; + _localLuminanceFactor = 1.0f; + _localLuminanceAddon = meanLuminance * _v0; + _maxInputValue = maxInputValue; + }; + inline void updateCompressionParameter(const float meanLuminance) + { + _localLuminanceFactor = 1; + _localLuminanceAddon = meanLuminance * _v0; + }; + inline float getV0CompressionParameter() + { + return _v0 / _maxInputValue; + }; + inline const cv::ocl::oclMat &getOutput() const + { + return _filterOutput; + }; + inline unsigned int getNBrows() + { + return _filterOutput.rows; + }; + inline unsigned int getNBcolumns() + { + return _filterOutput.cols; + }; + inline unsigned int getNBpixels() + { + return _filterOutput.size().area(); + }; + inline void normalizeGrayOutput_0_maxOutputValue(const float maxValue) + { + ocl::normalizeGrayOutput_0_maxOutputValue(_filterOutput, maxValue); + }; + inline void normalizeGrayOutputCentredSigmoide() + { + ocl::normalizeGrayOutputCentredSigmoide(0.0, 2.0, _filterOutput, _filterOutput); + }; + inline void centerReductImageLuminance() + { + ocl::centerReductImageLuminance(_filterOutput); + }; + inline float getMaxInputValue() + { + return this->_maxInputValue; + }; + inline void setMaxInputValue(const float newMaxInputValue) + { + this->_maxInputValue = newMaxInputValue; + }; + +protected: + cv::ocl::oclMat _filterOutput; + cv::ocl::oclMat _localBuffer; + + int _NBrows; + int _NBcols; + unsigned int _halfNBrows; + unsigned int _halfNBcolumns; + + std::valarray _filteringCoeficientsTable; + float _v0; + float _maxInputValue; + float _meanInputValue; + float _localLuminanceFactor; + float _localLuminanceAddon; + + float _a; + float _tau; + float _gain; + + void _spatiotemporalLPfilter(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &LPfilterOutput, const unsigned int coefTableOffset = 0); + float _squaringSpatiotemporalLPfilter(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame, const unsigned int filterIndex = 0); + void _spatiotemporalLPfilter_Irregular(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame, const unsigned int filterIndex = 0); + void _localSquaringSpatioTemporalLPfilter(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &LPfilterOutput, const unsigned int *integrationAreas, const unsigned int filterIndex = 0); + void _localLuminanceAdaptation(const cv::ocl::oclMat &inputFrame, const cv::ocl::oclMat &localLuminance, cv::ocl::oclMat &outputFrame, const bool updateLuminanceMean = true); + void _localLuminanceAdaptation(cv::ocl::oclMat &inputOutputFrame, const cv::ocl::oclMat &localLuminance); + void _localLuminanceAdaptationPosNegValues(const cv::ocl::oclMat &inputFrame, const cv::ocl::oclMat &localLuminance, float *outputFrame); + void _horizontalCausalFilter_addInput(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame); + void _horizontalAnticausalFilter(cv::ocl::oclMat &outputFrame); + void _verticalCausalFilter(cv::ocl::oclMat &outputFrame); + void _horizontalAnticausalFilter_Irregular(cv::ocl::oclMat &outputFrame, const cv::ocl::oclMat &spatialConstantBuffer); + void _verticalCausalFilter_Irregular(cv::ocl::oclMat &outputFrame, const cv::ocl::oclMat &spatialConstantBuffer); + void _verticalAnticausalFilter_multGain(cv::ocl::oclMat &outputFrame); +}; + +class MagnoRetinaFilter: public BasicRetinaFilter +{ +public: + MagnoRetinaFilter(const unsigned int NBrows, const unsigned int NBcolumns); + virtual ~MagnoRetinaFilter(); + void clearAllBuffers(); + void resize(const unsigned int NBrows, const unsigned int NBcolumns); + void setCoefficientsTable(const float parasolCells_beta, const float parasolCells_tau, const float parasolCells_k, const float amacrinCellsTemporalCutFrequency, const float localAdaptIntegration_tau, const float localAdaptIntegration_k); + + const cv::ocl::oclMat &runFilter(const cv::ocl::oclMat &OPL_ON, const cv::ocl::oclMat &OPL_OFF); + + inline const cv::ocl::oclMat &getMagnoON() const + { + return _magnoXOutputON; + }; + inline const cv::ocl::oclMat &getMagnoOFF() const + { + return _magnoXOutputOFF; + }; + inline const cv::ocl::oclMat &getMagnoYsaturated() const + { + return _magnoYsaturated; + }; + inline void normalizeGrayOutputNearZeroCentreredSigmoide() + { + ocl::normalizeGrayOutputNearZeroCentreredSigmoide(_magnoYOutput, _magnoYsaturated); + }; + inline float getTemporalConstant() + { + return this->_filteringCoeficientsTable[2]; + }; +private: + cv::ocl::oclMat _previousInput_ON; + cv::ocl::oclMat _previousInput_OFF; + cv::ocl::oclMat _amacrinCellsTempOutput_ON; + cv::ocl::oclMat _amacrinCellsTempOutput_OFF; + cv::ocl::oclMat _magnoXOutputON; + cv::ocl::oclMat _magnoXOutputOFF; + cv::ocl::oclMat _localProcessBufferON; + cv::ocl::oclMat _localProcessBufferOFF; + cv::ocl::oclMat _magnoYOutput; + cv::ocl::oclMat _magnoYsaturated; + + float _temporalCoefficient; + void _amacrineCellsComputing(const cv::ocl::oclMat &OPL_ON, const cv::ocl::oclMat &OPL_OFF); +}; + +class ParvoRetinaFilter: public BasicRetinaFilter +{ +public: + ParvoRetinaFilter(const unsigned int NBrows = 480, const unsigned int NBcolumns = 640); + virtual ~ParvoRetinaFilter(); + void resize(const unsigned int NBrows, const unsigned int NBcolumns); + void clearAllBuffers(); + void setOPLandParvoFiltersParameters(const float beta1, const float tau1, const float k1, const float beta2, const float tau2, const float k2); + + inline void setGanglionCellsLocalAdaptationLPfilterParameters(const float tau, const float k) + { + BasicRetinaFilter::setLPfilterParameters(0, tau, k, 2); + }; + const cv::ocl::oclMat &runFilter(const cv::ocl::oclMat &inputFrame, const bool useParvoOutput = true); + + inline const cv::ocl::oclMat &getPhotoreceptorsLPfilteringOutput() const + { + return _photoreceptorsOutput; + }; + + inline const cv::ocl::oclMat &getHorizontalCellsOutput() const + { + return _horizontalCellsOutput; + }; + + inline const cv::ocl::oclMat &getParvoON() const + { + return _parvocellularOutputON; + }; + + inline const cv::ocl::oclMat &getParvoOFF() const + { + return _parvocellularOutputOFF; + }; + + inline const cv::ocl::oclMat &getBipolarCellsON() const + { + return _bipolarCellsOutputON; + }; + + inline const cv::ocl::oclMat &getBipolarCellsOFF() const + { + return _bipolarCellsOutputOFF; + }; + + inline float getPhotoreceptorsTemporalConstant() + { + return this->_filteringCoeficientsTable[2]; + }; + + inline float getHcellsTemporalConstant() + { + return this->_filteringCoeficientsTable[5]; + }; +private: + cv::ocl::oclMat _photoreceptorsOutput; + cv::ocl::oclMat _horizontalCellsOutput; + cv::ocl::oclMat _parvocellularOutputON; + cv::ocl::oclMat _parvocellularOutputOFF; + cv::ocl::oclMat _bipolarCellsOutputON; + cv::ocl::oclMat _bipolarCellsOutputOFF; + cv::ocl::oclMat _localAdaptationOFF; + cv::ocl::oclMat _localAdaptationON; + cv::ocl::oclMat _parvocellularOutputONminusOFF; + void _OPL_OnOffWaysComputing(); +}; +class RetinaColor: public BasicRetinaFilter +{ +public: + RetinaColor(const unsigned int NBrows, const unsigned int NBcolumns, const int samplingMethod = RETINA_COLOR_DIAGONAL); + virtual ~RetinaColor(); + + void clearAllBuffers(); + void resize(const unsigned int NBrows, const unsigned int NBcolumns); + inline void runColorMultiplexing(const cv::ocl::oclMat &inputRGBFrame) + { + runColorMultiplexing(inputRGBFrame, _multiplexedFrame); + }; + void runColorMultiplexing(const cv::ocl::oclMat &demultiplexedInputFrame, cv::ocl::oclMat &multiplexedFrame); + void runColorDemultiplexing(const cv::ocl::oclMat &multiplexedColorFrame, const bool adaptiveFiltering = false, const float maxInputValue = 255.0); + + void setColorSaturation(const bool saturateColors = true, const float colorSaturationValue = 4.0) + { + _saturateColors = saturateColors; + _colorSaturationValue = colorSaturationValue; + }; + + void setChrominanceLPfilterParameters(const float beta, const float tau, const float k) + { + setLPfilterParameters(beta, tau, k); + }; + + bool applyKrauskopfLMS2Acr1cr2Transform(cv::ocl::oclMat &result); + bool applyLMS2LabTransform(cv::ocl::oclMat &result); + inline const cv::ocl::oclMat &getMultiplexedFrame() const + { + return _multiplexedFrame; + }; + + inline const cv::ocl::oclMat &getDemultiplexedColorFrame() const + { + return _demultiplexedColorFrame; + }; + + inline const cv::ocl::oclMat &getLuminance() const + { + return _luminance; + }; + inline const cv::ocl::oclMat &getChrominance() const + { + return _chrominance; + }; + void clipRGBOutput_0_maxInputValue(cv::ocl::oclMat &inputOutputBuffer, const float maxOutputValue = 255.0); + void normalizeRGBOutput_0_maxOutputValue(const float maxOutputValue = 255.0); + inline void setDemultiplexedColorFrame(const cv::ocl::oclMat &demultiplexedImage) + { + _demultiplexedColorFrame = demultiplexedImage; + }; +protected: + inline unsigned int bayerSampleOffset(unsigned int index) + { + return index + ((index / getNBcolumns()) % 2) * getNBpixels() + ((index % getNBcolumns()) % 2) * getNBpixels(); + } + inline Rect getROI(int idx) + { + return Rect(0, idx * _NBrows, _NBcols, _NBrows); + } + int _samplingMethod; + bool _saturateColors; + float _colorSaturationValue; + cv::ocl::oclMat _luminance; + cv::ocl::oclMat _multiplexedFrame; + cv::ocl::oclMat _RGBmosaic; + cv::ocl::oclMat _tempMultiplexedFrame; + cv::ocl::oclMat _demultiplexedTempBuffer; + cv::ocl::oclMat _demultiplexedColorFrame; + cv::ocl::oclMat _chrominance; + cv::ocl::oclMat _colorLocalDensity; + cv::ocl::oclMat _imageGradient; + + float _pR, _pG, _pB; + bool _objectInit; + + void _initColorSampling(); + void _adaptiveSpatialLPfilter(const cv::ocl::oclMat &inputFrame, const cv::ocl::oclMat &gradient, cv::ocl::oclMat &outputFrame); + void _adaptiveHorizontalCausalFilter_addInput(const cv::ocl::oclMat &inputFrame, const cv::ocl::oclMat &gradient, cv::ocl::oclMat &outputFrame); + void _adaptiveVerticalAnticausalFilter_multGain(const cv::ocl::oclMat &gradient, cv::ocl::oclMat &outputFrame); + void _computeGradient(const cv::ocl::oclMat &luminance, cv::ocl::oclMat &gradient); + void _normalizeOutputs_0_maxOutputValue(void); + void _applyImageColorSpaceConversion(const cv::ocl::oclMat &inputFrame, cv::ocl::oclMat &outputFrame, const float *transformTable); +}; +class RetinaFilter +{ +public: + RetinaFilter(const unsigned int sizeRows, const unsigned int sizeColumns, const bool colorMode = false, const int samplingMethod = RETINA_COLOR_BAYER, const bool useRetinaLogSampling = false, const double reductionFactor = 1.0, const double samplingStrenght = 10.0); + ~RetinaFilter(); + + void clearAllBuffers(); + void resize(const unsigned int NBrows, const unsigned int NBcolumns); + bool checkInput(const cv::ocl::oclMat &input, const bool colorMode); + bool runFilter(const cv::ocl::oclMat &imageInput, const bool useAdaptiveFiltering = true, const bool processRetinaParvoMagnoMapping = false, const bool useColorMode = false, const bool inputIsColorMultiplexed = false); + + void setGlobalParameters(const float OPLspatialResponse1 = 0.7, const float OPLtemporalresponse1 = 1, const float OPLassymetryGain = 0, const float OPLspatialResponse2 = 5, const float OPLtemporalresponse2 = 1, const float LPfilterSpatialResponse = 5, const float LPfilterGain = 0, const float LPfilterTemporalresponse = 0, const float MovingContoursExtractorCoefficient = 5, const bool normalizeParvoOutput_0_maxOutputValue = false, const bool normalizeMagnoOutput_0_maxOutputValue = false, const float maxOutputValue = 255.0, const float maxInputValue = 255.0, const float meanValue = 128.0); + + inline void setPhotoreceptorsLocalAdaptationSensitivity(const float V0CompressionParameter) + { + _photoreceptorsPrefilter.setV0CompressionParameter(1 - V0CompressionParameter); + _setInitPeriodCount(); + }; + + inline void setParvoGanglionCellsLocalAdaptationSensitivity(const float V0CompressionParameter) + { + _ParvoRetinaFilter.setV0CompressionParameter(V0CompressionParameter); + _setInitPeriodCount(); + }; + + inline void setGanglionCellsLocalAdaptationLPfilterParameters(const float spatialResponse, const float temporalResponse) + { + _ParvoRetinaFilter.setGanglionCellsLocalAdaptationLPfilterParameters(temporalResponse, spatialResponse); + _setInitPeriodCount(); + }; + + inline void setMagnoGanglionCellsLocalAdaptationSensitivity(const float V0CompressionParameter) + { + _MagnoRetinaFilter.setV0CompressionParameter(V0CompressionParameter); + _setInitPeriodCount(); + }; + + void setOPLandParvoParameters(const float beta1, const float tau1, const float k1, const float beta2, const float tau2, const float k2, const float V0CompressionParameter) + { + _ParvoRetinaFilter.setOPLandParvoFiltersParameters(beta1, tau1, k1, beta2, tau2, k2); + _ParvoRetinaFilter.setV0CompressionParameter(V0CompressionParameter); + _setInitPeriodCount(); + }; + + void setMagnoCoefficientsTable(const float parasolCells_beta, const float parasolCells_tau, const float parasolCells_k, const float amacrinCellsTemporalCutFrequency, const float V0CompressionParameter, const float localAdaptintegration_tau, const float localAdaptintegration_k) + { + _MagnoRetinaFilter.setCoefficientsTable(parasolCells_beta, parasolCells_tau, parasolCells_k, amacrinCellsTemporalCutFrequency, localAdaptintegration_tau, localAdaptintegration_k); + _MagnoRetinaFilter.setV0CompressionParameter(V0CompressionParameter); + _setInitPeriodCount(); + }; + + inline void activateNormalizeParvoOutput_0_maxOutputValue(const bool normalizeParvoOutput_0_maxOutputValue) + { + _normalizeParvoOutput_0_maxOutputValue = normalizeParvoOutput_0_maxOutputValue; + }; + + inline void activateNormalizeMagnoOutput_0_maxOutputValue(const bool normalizeMagnoOutput_0_maxOutputValue) + { + _normalizeMagnoOutput_0_maxOutputValue = normalizeMagnoOutput_0_maxOutputValue; + }; + + inline void setMaxOutputValue(const float maxOutputValue) + { + _maxOutputValue = maxOutputValue; + }; + + void setColorMode(const bool desiredColorMode) + { + _useColorMode = desiredColorMode; + }; + inline void setColorSaturation(const bool saturateColors = true, const float colorSaturationValue = 4.0) + { + _colorEngine.setColorSaturation(saturateColors, colorSaturationValue); + }; + inline const cv::ocl::oclMat &getLocalAdaptation() const + { + return _photoreceptorsPrefilter.getOutput(); + }; + inline const cv::ocl::oclMat &getPhotoreceptors() const + { + return _ParvoRetinaFilter.getPhotoreceptorsLPfilteringOutput(); + }; + + inline const cv::ocl::oclMat &getHorizontalCells() const + { + return _ParvoRetinaFilter.getHorizontalCellsOutput(); + }; + inline bool areContoursProcessed() + { + return _useParvoOutput; + }; + bool getParvoFoveaResponse(cv::ocl::oclMat &parvoFovealResponse); + inline void activateContoursProcessing(const bool useParvoOutput) + { + _useParvoOutput = useParvoOutput; + }; + + const cv::ocl::oclMat &getContours(); + + inline const cv::ocl::oclMat &getContoursON() const + { + return _ParvoRetinaFilter.getParvoON(); + }; + + inline const cv::ocl::oclMat &getContoursOFF() const + { + return _ParvoRetinaFilter.getParvoOFF(); + }; + + inline bool areMovingContoursProcessed() + { + return _useMagnoOutput; + }; + + inline void activateMovingContoursProcessing(const bool useMagnoOutput) + { + _useMagnoOutput = useMagnoOutput; + }; + + inline const cv::ocl::oclMat &getMovingContours() const + { + return _MagnoRetinaFilter.getOutput(); + }; + + inline const cv::ocl::oclMat &getMovingContoursSaturated() const + { + return _MagnoRetinaFilter.getMagnoYsaturated(); + }; + + inline const cv::ocl::oclMat &getMovingContoursON() const + { + return _MagnoRetinaFilter.getMagnoON(); + }; + + inline const cv::ocl::oclMat &getMovingContoursOFF() const + { + return _MagnoRetinaFilter.getMagnoOFF(); + }; + + inline const cv::ocl::oclMat &getRetinaParvoMagnoMappedOutput() const + { + return _retinaParvoMagnoMappedFrame; + }; + + inline const cv::ocl::oclMat &getParvoContoursChannel() const + { + return _colorEngine.getLuminance(); + }; + + inline const cv::ocl::oclMat &getParvoChrominance() const + { + return _colorEngine.getChrominance(); + }; + inline const cv::ocl::oclMat &getColorOutput() const + { + return _colorEngine.getDemultiplexedColorFrame(); + }; + + inline bool isColorMode() + { + return _useColorMode; + }; + bool getColorMode() + { + return _useColorMode; + }; + + inline bool isInitTransitionDone() + { + if (_ellapsedFramesSinceLastReset < _globalTemporalConstant) + { + return false; + } + return true; + }; + inline float getRetinaSamplingBackProjection(const float projectedRadiusLength) + { + return projectedRadiusLength; + }; + + inline unsigned int getInputNBrows() + { + return _photoreceptorsPrefilter.getNBrows(); + }; + + inline unsigned int getInputNBcolumns() + { + return _photoreceptorsPrefilter.getNBcolumns(); + }; + + inline unsigned int getInputNBpixels() + { + return _photoreceptorsPrefilter.getNBpixels(); + }; + + inline unsigned int getOutputNBrows() + { + return _photoreceptorsPrefilter.getNBrows(); + }; + + inline unsigned int getOutputNBcolumns() + { + return _photoreceptorsPrefilter.getNBcolumns(); + }; + + inline unsigned int getOutputNBpixels() + { + return _photoreceptorsPrefilter.getNBpixels(); + }; +private: + bool _useParvoOutput; + bool _useMagnoOutput; + + unsigned int _ellapsedFramesSinceLastReset; + unsigned int _globalTemporalConstant; + + cv::ocl::oclMat _retinaParvoMagnoMappedFrame; + BasicRetinaFilter _photoreceptorsPrefilter; + ParvoRetinaFilter _ParvoRetinaFilter; + MagnoRetinaFilter _MagnoRetinaFilter; + RetinaColor _colorEngine; + + bool _useMinimalMemoryForToneMappingONLY; + bool _normalizeParvoOutput_0_maxOutputValue; + bool _normalizeMagnoOutput_0_maxOutputValue; + float _maxOutputValue; + bool _useColorMode; + + void _setInitPeriodCount(); + void _processRetinaParvoMagnoMapping(); + void _runGrayToneMapping(const cv::ocl::oclMat &grayImageInput, cv::ocl::oclMat &grayImageOutput , const float PhotoreceptorsCompression = 0.6, const float ganglionCellsCompression = 0.6); +}; + +} /* namespace ocl */ +} /* namespace bioinspired */ +} /* namespace cv */ + +#endif /* HAVE_OPENCV_OCL */ +#endif /* __OCL_RETINA_HPP__ */ diff --git a/modules/bioinspired/test/test_retina_ocl.cpp b/modules/bioinspired/test/test_retina_ocl.cpp new file mode 100644 index 0000000..ea40bf6 --- /dev/null +++ b/modules/bioinspired/test/test_retina_ocl.cpp @@ -0,0 +1,139 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2013, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2013, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Peng Xiao, pengxiao@multicorewareinc.com +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other oclMaterials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors as is and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +#include "test_precomp.hpp" +#include "opencv2/opencv_modules.hpp" +#include "opencv2/bioinspired.hpp" +#include "opencv2/imgproc.hpp" +#include "opencv2/highgui.hpp" +#include "opencv2/ocl.hpp" + +#if defined(HAVE_OPENCV_OCL) && defined(HAVE_OPENCL) + +#define RETINA_ITERATIONS 5 + +static double checkNear(const cv::Mat &m1, const cv::Mat &m2) +{ + return cv::norm(m1, m2, cv::NORM_INF); +} + +#define PARAM_TEST_CASE(name, ...) struct name : testing::TestWithParam< std::tr1::tuple< __VA_ARGS__ > > +#define GET_PARAM(k) std::tr1::get< k >(GetParam()) + + +PARAM_TEST_CASE(Retina_OCL, bool, int, bool, double, double) +{ + bool colorMode; + int colorSamplingMethod; + bool useLogSampling; + double reductionFactor; + double samplingStrength; + + std::vector infos; + + virtual void SetUp() + { + colorMode = GET_PARAM(0); + colorSamplingMethod = GET_PARAM(1); + useLogSampling = GET_PARAM(2); + reductionFactor = GET_PARAM(3); + samplingStrength = GET_PARAM(4); + + cv::ocl::getDevice(infos); + std::cout << "Device name:" << infos[0].DeviceName[0] << std::endl; + } +}; + +TEST_P(Retina_OCL, Accuracy) +{ + using namespace cv; + Mat input = imread(cvtest::TS::ptr()->get_data_path() + "shared/lena.png", colorMode); + CV_Assert(!input.empty()); + ocl::oclMat ocl_input(input); + + Ptr ocl_retina = bioinspired::createRetina_OCL( + input.size(), + colorMode, + colorSamplingMethod, + useLogSampling, + reductionFactor, + samplingStrength); + + Ptr gold_retina = bioinspired::createRetina( + input.size(), + colorMode, + colorSamplingMethod, + useLogSampling, + reductionFactor, + samplingStrength); + + Mat gold_parvo; + Mat gold_magno; + ocl::oclMat ocl_parvo; + ocl::oclMat ocl_magno; + + for(int i = 0; i < RETINA_ITERATIONS; i ++) + { + ocl_retina->run(ocl_input); + gold_retina->run(input); + + gold_retina->getParvo(gold_parvo); + gold_retina->getMagno(gold_magno); + + ocl_retina->getParvo(ocl_parvo); + ocl_retina->getMagno(ocl_magno); + + EXPECT_LE(checkNear(gold_parvo, (Mat)ocl_parvo), 1.0); + EXPECT_LE(checkNear(gold_magno, (Mat)ocl_magno), 1.0); + } +} + +INSTANTIATE_TEST_CASE_P(Contrib, Retina_OCL, testing::Combine( + testing::Values(false, true), + testing::Values((int)cv::bioinspired::RETINA_COLOR_BAYER), + testing::Values(false/*,true*/), + testing::Values(1.0, 0.5), + testing::Values(10.0, 5.0))); +#endif -- 2.7.4