Add OpenCL accelerated implementation for Retina.
authorpeng xiao <hisenxpress@gmail.com>
Tue, 13 Aug 2013 08:27:12 +0000 (16:27 +0800)
committerpeng xiao <hisenxpress@gmail.com>
Tue, 13 Aug 2013 08:27:12 +0000 (16:27 +0800)
modules/bioinspired/CMakeLists.txt
modules/bioinspired/include/opencv2/bioinspired/retina.hpp
modules/bioinspired/src/opencl/retina_kernel.cl [new file with mode: 0644]
modules/bioinspired/src/precomp.hpp
modules/bioinspired/src/retina_ocl.cpp [new file with mode: 0644]
modules/bioinspired/src/retina_ocl.hpp [new file with mode: 0644]
modules/bioinspired/test/test_retina_ocl.cpp [new file with mode: 0644]

index a27ad73..b0f152c 100644 (file)
@@ -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)
index 4b4c19f..b4fda70 100644 (file)
@@ -304,7 +304,8 @@ public:
 CV_EXPORTS Ptr<Retina> createRetina(Size inputSize);
 CV_EXPORTS Ptr<Retina> 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<Retina> createRetina_OCL(Size inputSize);
+CV_EXPORTS Ptr<Retina> 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 (file)
index 0000000..ed5f831
--- /dev/null
@@ -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;
+}
index 066e15b..541b970 100644 (file)
 #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 <valarray>
+
+#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 (file)
index 0000000..ba98da2
--- /dev/null
@@ -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 <iostream>
+#include <sstream>
+
+#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<float>
+    // 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<oclMat> channel_splits(numOfSplits);
+        for(int i = 0; i < static_cast<int>(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<float>(ocl::sum(inputFrame)[0]) / getNBpixels();
+        updateCompressionParameter(meanLuminance);
+    }
+    int elements_per_row = static_cast<int>(inputFrame.step / inputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(inputFrame.step / inputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<float>(max_val - min_val);
+    float offset = - static_cast<float>(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<type>::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<std::pair<size_t, const void *> > args;
+    size_t globalSize[] = {in.cols, out.rows, 1};
+    size_t localSize[]  = {16, 16, 1};
+
+    int elements_per_row = static_cast<int>(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<std::pair<size_t, const void *> > args;
+    size_t globalSize[] = {inputPicture.cols, inputPicture.rows, 1};
+    size_t localSize[]  = {16, 16, 1};
+
+    int elements_per_row = static_cast<int>(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<std::pair<size_t, const void *> > args;
+    size_t globalSize[] = {inputoutput.cols, inputoutput.rows, 1};
+    size_t localSize[]  = {16, 16, 1};
+
+    float f_mean = static_cast<float>(mean[0]);
+    float f_stddev = static_cast<float>(stddev[0]);
+    int elements_per_row = static_cast<int>(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<int>(_photoreceptorsOutput.step / _photoreceptorsOutput.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(OPL_ON.step / OPL_ON.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(input.step / input.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<float>();
+    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<int>(input.step / input.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(ocl_luma.step / ocl_luma.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(colorDemultiplex.step / colorDemultiplex.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(input.step / input.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(multiplexedFrame.step / multiplexedFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(inputOutputBuffer.step / inputOutputBuffer.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(inputFrame.step / inputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(outputFrame.step / outputFrame.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > args;
+    size_t globalSize[] = {_NBcols, 1, 1};
+    size_t localSize[]  = {256, 1, 1};
+
+    int gradOffset = gradient.offset + static_cast<int>(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<int>(luminance.step / luminance.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<int>(inputTarget->getNBrows())
+                 || input.rows == static_cast<int>(inputTarget->getNBrows()) * 3
+                 || input.rows == static_cast<int>(inputTarget->getNBrows()) * 4)
+                && input.cols == static_cast<int>(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<int>(parvo.step / parvo.elemSize());
+
+    Context * ctx = Context::getContext();
+    std::vector<std::pair<size_t, const void *> > 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<Retina> createRetina_OCL(Size getInputSize){ return new ocl::RetinaOCLImpl(getInputSize); }
+Ptr<Retina> 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 (file)
index 0000000..69a833c
--- /dev/null
@@ -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 <float>_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 (file)
index 0000000..ea40bf6
--- /dev/null
@@ -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<cv::ocl::Info> 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<bioinspired::Retina> ocl_retina = bioinspired::createRetina_OCL(
+        input.size(),
+        colorMode,
+        colorSamplingMethod,
+        useLogSampling,
+        reductionFactor,
+        samplingStrength);
+
+    Ptr<bioinspired::Retina> 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