Add ocl version of boxFilter to the module img_proc
This commit is contained in:
		
							
								
								
									
										379
									
								
								modules/imgproc/src/opencl/boxFilter.cl
									
									
									
									
									
										Normal file
									
								
							
							
						
						
									
										379
									
								
								modules/imgproc/src/opencl/boxFilter.cl
									
									
									
									
									
										Normal file
									
								
							@@ -0,0 +1,379 @@
 | 
			
		||||
/*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, Advanced Micro Devices, Inc., all rights reserved.
 | 
			
		||||
// Third party copyrights are property of their respective owners.
 | 
			
		||||
//
 | 
			
		||||
// 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 materials 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*/
 | 
			
		||||
 | 
			
		||||
///////////////////////////////////////////////////////////////////////////////////////////////////
 | 
			
		||||
/////////////////////////////////Macro for border type////////////////////////////////////////////
 | 
			
		||||
/////////////////////////////////////////////////////////////////////////////////////////////////
 | 
			
		||||
#ifdef BORDER_REPLICATE
 | 
			
		||||
//BORDER_REPLICATE:     aaaaaa|abcdefgh|hhhhhhh
 | 
			
		||||
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? (l_edge)   : (i))
 | 
			
		||||
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? (r_edge)-1 : (addr))
 | 
			
		||||
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? (t_edge)   :(i))
 | 
			
		||||
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? (b_edge)-1 :(addr))
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#ifdef BORDER_REFLECT
 | 
			
		||||
//BORDER_REFLECT:       fedcba|abcdefgh|hgfedcb
 | 
			
		||||
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? -(i)-1               : (i))
 | 
			
		||||
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr))
 | 
			
		||||
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? -(i)-1 : (i))
 | 
			
		||||
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? -(i)-1+((b_edge)<<1) : (addr))
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#ifdef BORDER_REFLECT_101
 | 
			
		||||
//BORDER_REFLECT_101:   gfedcb|abcdefgh|gfedcba
 | 
			
		||||
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? -(i)                 : (i))
 | 
			
		||||
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr))
 | 
			
		||||
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? -(i)                 : (i))
 | 
			
		||||
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? -(i)-2+((b_edge)<<1) : (addr))
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
//blur function does not support BORDER_WRAP
 | 
			
		||||
#ifdef BORDER_WRAP
 | 
			
		||||
//BORDER_WRAP:          cdefgh|abcdefgh|abcdefg
 | 
			
		||||
#define ADDR_L(i, l_edge, r_edge)  ((i) <  (l_edge) ? (i)+(r_edge) : (i))
 | 
			
		||||
#define ADDR_R(i, r_edge, addr)    ((i) >= (r_edge) ? (i)-(r_edge) : (addr))
 | 
			
		||||
#define ADDR_H(i, t_edge, b_edge)  ((i) <  (t_edge) ? (i)+(b_edge) : (i))
 | 
			
		||||
#define ADDR_B(i, b_edge, addr)    ((i) >= (b_edge) ? (i)-(b_edge) : (addr))
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#ifdef EXTRA_EXTRAPOLATION // border > src image size
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
// None
 | 
			
		||||
#elif defined BORDER_REPLICATE
 | 
			
		||||
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 | 
			
		||||
    { \
 | 
			
		||||
        x = max(min(x, maxX - 1), minX); \
 | 
			
		||||
        y = max(min(y, maxY - 1), minY); \
 | 
			
		||||
    }
 | 
			
		||||
#elif defined BORDER_WRAP
 | 
			
		||||
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 | 
			
		||||
    { \
 | 
			
		||||
        if (x < minX) \
 | 
			
		||||
            x -= ((x - maxX + 1) / maxX) * maxX; \
 | 
			
		||||
        if (x >= maxX) \
 | 
			
		||||
            x %= maxX; \
 | 
			
		||||
        if (y < minY) \
 | 
			
		||||
            y -= ((y - maxY + 1) / maxY) * maxY; \
 | 
			
		||||
        if (y >= maxY) \
 | 
			
		||||
            y %= maxY; \
 | 
			
		||||
    }
 | 
			
		||||
#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101)
 | 
			
		||||
#define EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, delta) \
 | 
			
		||||
    { \
 | 
			
		||||
        if (maxX - minX == 1) \
 | 
			
		||||
            x = minX; \
 | 
			
		||||
        else \
 | 
			
		||||
            do \
 | 
			
		||||
            { \
 | 
			
		||||
                if (x < minX) \
 | 
			
		||||
                    x = -(x - minX) - 1 + delta; \
 | 
			
		||||
                else \
 | 
			
		||||
                    x = maxX - 1 - (x - maxX) - delta; \
 | 
			
		||||
            } \
 | 
			
		||||
            while (x >= maxX || x < minX); \
 | 
			
		||||
        \
 | 
			
		||||
        if (maxY - minY == 1) \
 | 
			
		||||
            y = minY; \
 | 
			
		||||
        else \
 | 
			
		||||
            do \
 | 
			
		||||
            { \
 | 
			
		||||
                if (y < minY) \
 | 
			
		||||
                    y = -(y - minY) - 1 + delta; \
 | 
			
		||||
                else \
 | 
			
		||||
                    y = maxY - 1 - (y - maxY) - delta; \
 | 
			
		||||
            } \
 | 
			
		||||
            while (y >= maxY || y < minY); \
 | 
			
		||||
    }
 | 
			
		||||
#ifdef BORDER_REFLECT
 | 
			
		||||
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 0)
 | 
			
		||||
#elif defined(BORDER_REFLECT_101)
 | 
			
		||||
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) EXTRAPOLATE_(x, y, minX, minY, maxX, maxY, 1)
 | 
			
		||||
#endif
 | 
			
		||||
#else
 | 
			
		||||
#error No extrapolation method
 | 
			
		||||
#endif
 | 
			
		||||
#else
 | 
			
		||||
#define EXTRAPOLATE(x, y, minX, minY, maxX, maxY) \
 | 
			
		||||
    { \
 | 
			
		||||
        int _row = y - minY, _col = x - minX; \
 | 
			
		||||
        _row = ADDR_H(_row, 0, maxY - minY); \
 | 
			
		||||
        _row = ADDR_B(_row, maxY - minY, _row); \
 | 
			
		||||
        y = _row + minY; \
 | 
			
		||||
        \
 | 
			
		||||
        _col = ADDR_L(_col, 0, maxX - minX); \
 | 
			
		||||
        _col = ADDR_R(_col, maxX - minX, _col); \
 | 
			
		||||
        x = _col + minX; \
 | 
			
		||||
    }
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#if USE_DOUBLE
 | 
			
		||||
#ifdef cl_amd_fp64
 | 
			
		||||
#pragma OPENCL EXTENSION cl_amd_fp64:enable
 | 
			
		||||
#elif defined (cl_khr_fp64)
 | 
			
		||||
#pragma OPENCL EXTENSION cl_khr_fp64:enable
 | 
			
		||||
#endif
 | 
			
		||||
#define FPTYPE double
 | 
			
		||||
#define CONVERT_TO_FPTYPE CAT(convert_double, VEC_SIZE)
 | 
			
		||||
#else
 | 
			
		||||
#define FPTYPE float
 | 
			
		||||
#define CONVERT_TO_FPTYPE CAT(convert_float, VEC_SIZE)
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#if DATA_DEPTH == 0
 | 
			
		||||
#define BASE_TYPE uchar
 | 
			
		||||
#elif DATA_DEPTH == 1
 | 
			
		||||
#define BASE_TYPE char
 | 
			
		||||
#elif DATA_DEPTH == 2
 | 
			
		||||
#define BASE_TYPE ushort
 | 
			
		||||
#elif DATA_DEPTH == 3
 | 
			
		||||
#define BASE_TYPE short
 | 
			
		||||
#elif DATA_DEPTH == 4
 | 
			
		||||
#define BASE_TYPE int
 | 
			
		||||
#elif DATA_DEPTH == 5
 | 
			
		||||
#define BASE_TYPE float
 | 
			
		||||
#elif DATA_DEPTH == 6
 | 
			
		||||
#define BASE_TYPE double
 | 
			
		||||
#else
 | 
			
		||||
#error data_depth
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#define __CAT(x, y) x##y
 | 
			
		||||
#define CAT(x, y) __CAT(x, y)
 | 
			
		||||
 | 
			
		||||
#define uchar1 uchar
 | 
			
		||||
#define char1 char
 | 
			
		||||
#define ushort1 ushort
 | 
			
		||||
#define short1 short
 | 
			
		||||
#define int1 int
 | 
			
		||||
#define float1 float
 | 
			
		||||
#define double1 double
 | 
			
		||||
 | 
			
		||||
#define convert_uchar1_sat_rte convert_uchar_sat_rte
 | 
			
		||||
#define convert_char1_sat_rte convert_char_sat_rte
 | 
			
		||||
#define convert_ushort1_sat_rte convert_ushort_sat_rte
 | 
			
		||||
#define convert_short1_sat_rte convert_short_sat_rte
 | 
			
		||||
#define convert_int1_sat_rte convert_int_sat_rte
 | 
			
		||||
#define convert_float1
 | 
			
		||||
#define convert_double1
 | 
			
		||||
 | 
			
		||||
#if DATA_DEPTH == 5 || DATA_DEPTH == 6
 | 
			
		||||
#define CONVERT_TO_TYPE CAT(CAT(convert_, BASE_TYPE), VEC_SIZE)
 | 
			
		||||
#else
 | 
			
		||||
#define CONVERT_TO_TYPE CAT(CAT(CAT(convert_, BASE_TYPE), VEC_SIZE), _sat_rte)
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
#define VEC_SIZE DATA_CHAN
 | 
			
		||||
 | 
			
		||||
#define VEC_TYPE CAT(BASE_TYPE, VEC_SIZE)
 | 
			
		||||
#define TYPE VEC_TYPE
 | 
			
		||||
 | 
			
		||||
#define SCALAR_TYPE CAT(FPTYPE, VEC_SIZE)
 | 
			
		||||
 | 
			
		||||
#define INTERMEDIATE_TYPE CAT(FPTYPE, VEC_SIZE)
 | 
			
		||||
 | 
			
		||||
#define TYPE_SIZE   (VEC_SIZE*sizeof(BASE_TYPE))
 | 
			
		||||
 | 
			
		||||
struct RectCoords
 | 
			
		||||
{
 | 
			
		||||
    int x1, y1, x2, y2;
 | 
			
		||||
};
 | 
			
		||||
 | 
			
		||||
//#define DEBUG
 | 
			
		||||
#ifdef DEBUG
 | 
			
		||||
#define DEBUG_ONLY(x) x
 | 
			
		||||
#define ASSERT(condition) do { if (!(condition)) { printf("BUG in boxFilter kernel (global=%d,%d): " #condition "\n", get_global_id(0), get_global_id(1)); } } while (0)
 | 
			
		||||
#else
 | 
			
		||||
#define DEBUG_ONLY(x)
 | 
			
		||||
#define ASSERT(condition)
 | 
			
		||||
#endif
 | 
			
		||||
 | 
			
		||||
 | 
			
		||||
inline INTERMEDIATE_TYPE readSrcPixel(int2 pos, __global const uchar* srcptr, int srcstep, int srcoffset, const struct RectCoords srcCoords
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
               , SCALAR_TYPE borderValue
 | 
			
		||||
#endif
 | 
			
		||||
    )
 | 
			
		||||
{
 | 
			
		||||
#ifdef BORDER_ISOLATED
 | 
			
		||||
    if(pos.x >= srcCoords.x1 && pos.y >= srcCoords.y1 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
 | 
			
		||||
#else
 | 
			
		||||
    if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
 | 
			
		||||
#endif
 | 
			
		||||
    {
 | 
			
		||||
        __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * TYPE_SIZE/*sizeof(TYPE)*/);
 | 
			
		||||
        return CONVERT_TO_FPTYPE(*ptr);
 | 
			
		||||
    }
 | 
			
		||||
    else
 | 
			
		||||
    {
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
        return borderValue;
 | 
			
		||||
#else
 | 
			
		||||
        int selected_col = pos.x;
 | 
			
		||||
        int selected_row = pos.y;
 | 
			
		||||
 | 
			
		||||
        EXTRAPOLATE(selected_col, selected_row,
 | 
			
		||||
#ifdef BORDER_ISOLATED
 | 
			
		||||
                srcCoords.x1, srcCoords.y1,
 | 
			
		||||
#else
 | 
			
		||||
                0, 0,
 | 
			
		||||
#endif
 | 
			
		||||
                srcCoords.x2, srcCoords.y2
 | 
			
		||||
         );
 | 
			
		||||
 | 
			
		||||
        // debug border mapping
 | 
			
		||||
        //printf("pos=%d,%d --> %d, %d\n", pos.x, pos.y, selected_col, selected_row);
 | 
			
		||||
 | 
			
		||||
        pos = (int2)(selected_col, selected_row);
 | 
			
		||||
        if(pos.x >= 0 && pos.y >= 0 && pos.x < srcCoords.x2 && pos.y < srcCoords.y2)
 | 
			
		||||
        {
 | 
			
		||||
            __global TYPE* ptr = (__global TYPE*)(srcptr + pos.y * srcstep + srcoffset + pos.x * TYPE_SIZE/*sizeof(TYPE)*/);
 | 
			
		||||
            return CONVERT_TO_FPTYPE(*ptr);
 | 
			
		||||
        }
 | 
			
		||||
        else
 | 
			
		||||
        {
 | 
			
		||||
            // for debug only
 | 
			
		||||
            DEBUG_ONLY(printf("BUG in boxFilter kernel\n"));
 | 
			
		||||
            return (FPTYPE)(0.0f);
 | 
			
		||||
        }
 | 
			
		||||
#endif
 | 
			
		||||
    }
 | 
			
		||||
}
 | 
			
		||||
 | 
			
		||||
// INPUT PARAMETER: BLOCK_SIZE_Y (via defines)
 | 
			
		||||
 | 
			
		||||
__kernel
 | 
			
		||||
__attribute__((reqd_work_group_size(LOCAL_SIZE, 1, 1)))
 | 
			
		||||
void boxFilter(__global const uchar* srcptr, int srcstep, int srcoffset,
 | 
			
		||||
               __global uchar* dstptr, int dststep, int dstoffset,
 | 
			
		||||
               int rows, int cols,
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
               SCALAR_TYPE borderValue,
 | 
			
		||||
#endif
 | 
			
		||||
               FPTYPE alpha
 | 
			
		||||
               )
 | 
			
		||||
{
 | 
			
		||||
    const struct RectCoords srcCoords = {0, 0, cols, rows}; // for non-isolated border: offsetX, offsetY, wholeX, wholeY
 | 
			
		||||
    const struct RectCoords dstCoords = {0, 0, cols, rows};
 | 
			
		||||
 | 
			
		||||
    const int x = get_local_id(0) + (LOCAL_SIZE - (KERNEL_SIZE_X - 1)) * get_group_id(0) - ANCHOR_X;
 | 
			
		||||
    const int y = get_global_id(1) * BLOCK_SIZE_Y;
 | 
			
		||||
 | 
			
		||||
    const int local_id = get_local_id(0);
 | 
			
		||||
 | 
			
		||||
    INTERMEDIATE_TYPE data[KERNEL_SIZE_Y];
 | 
			
		||||
    __local INTERMEDIATE_TYPE sumOfCols[LOCAL_SIZE];
 | 
			
		||||
 | 
			
		||||
    int2 srcPos = (int2)(srcCoords.x1 + x, srcCoords.y1 + y - ANCHOR_Y);
 | 
			
		||||
    for(int sy = 0; sy < KERNEL_SIZE_Y; sy++, srcPos.y++)
 | 
			
		||||
    {
 | 
			
		||||
        data[sy] = readSrcPixel(srcPos, srcptr, srcstep, srcoffset, srcCoords
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
                , borderValue
 | 
			
		||||
#endif
 | 
			
		||||
                );
 | 
			
		||||
    }
 | 
			
		||||
 | 
			
		||||
    INTERMEDIATE_TYPE tmp_sum = 0;
 | 
			
		||||
    for(int sy = 0; sy < KERNEL_SIZE_Y; sy++)
 | 
			
		||||
    {
 | 
			
		||||
        tmp_sum += (data[sy]);
 | 
			
		||||
    }
 | 
			
		||||
 | 
			
		||||
    sumOfCols[local_id] = tmp_sum;
 | 
			
		||||
    barrier(CLK_LOCAL_MEM_FENCE);
 | 
			
		||||
 | 
			
		||||
    int2 pos = (int2)(dstCoords.x1 + x, dstCoords.y1 + y);
 | 
			
		||||
    __global TYPE* dstPtr = (__global TYPE*)(dstptr + pos.y * dststep + dstoffset + pos.x * TYPE_SIZE/*sizeof(TYPE)*/); // Pointer can be out of bounds!
 | 
			
		||||
 | 
			
		||||
    int sy_index = 0; // current index in data[] array
 | 
			
		||||
    int stepsY = min(dstCoords.y2 - pos.y, BLOCK_SIZE_Y);
 | 
			
		||||
    ASSERT(stepsY > 0);
 | 
			
		||||
    for (; ;)
 | 
			
		||||
    {
 | 
			
		||||
        ASSERT(pos.y < dstCoords.y2);
 | 
			
		||||
 | 
			
		||||
        if(local_id >= ANCHOR_X && local_id < LOCAL_SIZE - (KERNEL_SIZE_X - 1 - ANCHOR_X) &&
 | 
			
		||||
            pos.x >= dstCoords.x1 && pos.x < dstCoords.x2)
 | 
			
		||||
        {
 | 
			
		||||
            ASSERT(pos.y >= dstCoords.y1 && pos.y < dstCoords.y2);
 | 
			
		||||
 | 
			
		||||
            INTERMEDIATE_TYPE total_sum = 0;
 | 
			
		||||
#pragma unroll
 | 
			
		||||
            for (int sx = 0; sx < KERNEL_SIZE_X; sx++)
 | 
			
		||||
            {
 | 
			
		||||
                total_sum += sumOfCols[local_id + sx - ANCHOR_X];
 | 
			
		||||
            }
 | 
			
		||||
            *dstPtr = CONVERT_TO_TYPE(((INTERMEDIATE_TYPE)alpha) * total_sum);
 | 
			
		||||
        }
 | 
			
		||||
 | 
			
		||||
#if BLOCK_SIZE_Y == 1
 | 
			
		||||
        break;
 | 
			
		||||
#else
 | 
			
		||||
        if (--stepsY == 0)
 | 
			
		||||
            break;
 | 
			
		||||
 | 
			
		||||
        barrier(CLK_LOCAL_MEM_FENCE);
 | 
			
		||||
 | 
			
		||||
        tmp_sum = sumOfCols[local_id]; // TODO FIX IT: workaround for BUG in OpenCL compiler
 | 
			
		||||
        // only works with scalars: ASSERT(fabs(tmp_sum - sumOfCols[local_id]) < (INTERMEDIATE_TYPE)1e-6);
 | 
			
		||||
        tmp_sum -= data[sy_index];
 | 
			
		||||
 | 
			
		||||
        data[sy_index] = readSrcPixel(srcPos, srcptr, srcstep, srcoffset, srcCoords
 | 
			
		||||
#ifdef BORDER_CONSTANT
 | 
			
		||||
                , borderValue
 | 
			
		||||
#endif
 | 
			
		||||
                );
 | 
			
		||||
        srcPos.y++;
 | 
			
		||||
 | 
			
		||||
        tmp_sum += data[sy_index];
 | 
			
		||||
        sumOfCols[local_id] = tmp_sum;
 | 
			
		||||
 | 
			
		||||
        sy_index = (sy_index + 1 < KERNEL_SIZE_Y) ? sy_index + 1 : 0;
 | 
			
		||||
 | 
			
		||||
        barrier(CLK_LOCAL_MEM_FENCE);
 | 
			
		||||
 | 
			
		||||
        // next line
 | 
			
		||||
        DEBUG_ONLY(pos.y++);
 | 
			
		||||
        dstPtr = (__global TYPE*)((__global char*)dstPtr + dststep); // Pointer can be out of bounds!
 | 
			
		||||
#endif // BLOCK_SIZE_Y == 1
 | 
			
		||||
    }
 | 
			
		||||
}
 | 
			
		||||
@@ -41,6 +41,7 @@
 | 
			
		||||
//M*/
 | 
			
		||||
 | 
			
		||||
#include "precomp.hpp"
 | 
			
		||||
#include "opencl_kernels.hpp"
 | 
			
		||||
 | 
			
		||||
/*
 | 
			
		||||
 * This file includes the code, contributed by Simon Perreault
 | 
			
		||||
@@ -610,9 +611,137 @@ template<> struct ColumnSum<int, ushort> : public BaseColumnFilter
 | 
			
		||||
    std::vector<int> sum;
 | 
			
		||||
};
 | 
			
		||||
 | 
			
		||||
#define DIVUP(total, grain) ((total + grain - 1) / (grain))
 | 
			
		||||
 | 
			
		||||
static bool ocl_boxFilter( InputArray _src, OutputArray _dst, int ddepth,
 | 
			
		||||
                           Size ksize, Point anchor, int borderType )
 | 
			
		||||
{
 | 
			
		||||
    int type = _src.type();
 | 
			
		||||
    int cn = CV_MAT_CN(type);
 | 
			
		||||
    if ((1 != cn) && (2 != cn) && (4 != cn))
 | 
			
		||||
        return false;//TODO
 | 
			
		||||
 | 
			
		||||
    int sdepth = CV_MAT_DEPTH(type);
 | 
			
		||||
    if( anchor.x < 0 )
 | 
			
		||||
        anchor.x = ksize.width / 2;
 | 
			
		||||
    if( anchor.y < 0 )
 | 
			
		||||
        anchor.y = ksize.height / 2;
 | 
			
		||||
    if( ddepth < 0 )
 | 
			
		||||
        ddepth = sdepth;
 | 
			
		||||
 | 
			
		||||
    ocl::Kernel kernel;
 | 
			
		||||
 | 
			
		||||
    //Normalize the result by default
 | 
			
		||||
    float alpha = 1.0f / (ksize.height * ksize.width);
 | 
			
		||||
    bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0;
 | 
			
		||||
    bool useDouble = (CV_64F == sdepth);
 | 
			
		||||
    const cv::ocl::Device &device = cv::ocl::Device::getDefault();
 | 
			
		||||
    int doubleFPConfig = device.doubleFPConfig();
 | 
			
		||||
    if (useDouble && (0 == doubleFPConfig))
 | 
			
		||||
        return false;// may be we have to check is  (0 != (CL_FP_SOFT_FLOAT & doubleFPConfig)) ?
 | 
			
		||||
 | 
			
		||||
    const char* btype = NULL;
 | 
			
		||||
    switch (borderType & ~BORDER_ISOLATED)
 | 
			
		||||
    {
 | 
			
		||||
    case BORDER_CONSTANT:
 | 
			
		||||
        btype = "BORDER_CONSTANT";
 | 
			
		||||
        break;
 | 
			
		||||
    case BORDER_REPLICATE:
 | 
			
		||||
        btype = "BORDER_REPLICATE";
 | 
			
		||||
        break;
 | 
			
		||||
    case BORDER_REFLECT:
 | 
			
		||||
        btype = "BORDER_REFLECT";
 | 
			
		||||
        break;
 | 
			
		||||
    case BORDER_WRAP:
 | 
			
		||||
        //CV_Error(CV_StsUnsupportedFormat, "BORDER_WRAP is not supported!");
 | 
			
		||||
        return false;
 | 
			
		||||
    case BORDER_REFLECT101:
 | 
			
		||||
        btype = "BORDER_REFLECT_101";
 | 
			
		||||
        break;
 | 
			
		||||
    }
 | 
			
		||||
 | 
			
		||||
    cv::Size sz = _src.size();
 | 
			
		||||
 | 
			
		||||
    size_t globalsize[2] = {sz.width, sz.height};
 | 
			
		||||
    size_t localsize[2] = {0, 1};
 | 
			
		||||
 | 
			
		||||
    size_t maxWorkItemSizes[32]; device.maxWorkItemSizes(maxWorkItemSizes);
 | 
			
		||||
    size_t tryWorkItems = maxWorkItemSizes[0];
 | 
			
		||||
    do {
 | 
			
		||||
        size_t BLOCK_SIZE = tryWorkItems;
 | 
			
		||||
        while (BLOCK_SIZE > 32 && BLOCK_SIZE >= (size_t)ksize.width * 2 && BLOCK_SIZE > (size_t)sz.width * 2)
 | 
			
		||||
            BLOCK_SIZE /= 2;
 | 
			
		||||
        size_t BLOCK_SIZE_Y = 8; // TODO Check heuristic value on devices
 | 
			
		||||
        while (BLOCK_SIZE_Y < BLOCK_SIZE / 8 && BLOCK_SIZE_Y * device.maxComputeUnits() * 32 < (size_t)sz.height)
 | 
			
		||||
            BLOCK_SIZE_Y *= 2;
 | 
			
		||||
 | 
			
		||||
        if ((size_t)ksize.width > BLOCK_SIZE)
 | 
			
		||||
            return false;
 | 
			
		||||
 | 
			
		||||
        int requiredTop = anchor.y;
 | 
			
		||||
        int requiredLeft = (int)BLOCK_SIZE; // not this: anchor.x;
 | 
			
		||||
        int requiredBottom = ksize.height - 1 - anchor.y;
 | 
			
		||||
        int requiredRight = (int)BLOCK_SIZE; // not this: ksize.width - 1 - anchor.x;
 | 
			
		||||
        int h = sz.height;
 | 
			
		||||
        int w = sz.width;
 | 
			
		||||
        bool extra_extrapolation = h < requiredTop || h < requiredBottom || w < requiredLeft || w < requiredRight;
 | 
			
		||||
 | 
			
		||||
        if ((w < ksize.width) || (h < ksize.height))
 | 
			
		||||
            return false;
 | 
			
		||||
 | 
			
		||||
        char build_options[1024];
 | 
			
		||||
        sprintf(build_options, "-D LOCAL_SIZE=%d -D BLOCK_SIZE_Y=%d -D DATA_DEPTH=%d -D DATA_CHAN=%d -D USE_DOUBLE=%d -D ANCHOR_X=%d -D ANCHOR_Y=%d -D KERNEL_SIZE_X=%d -D KERNEL_SIZE_Y=%d -D %s -D %s -D %s",
 | 
			
		||||
                (int)BLOCK_SIZE, (int)BLOCK_SIZE_Y,
 | 
			
		||||
                sdepth, cn, useDouble ? 1 : 0,
 | 
			
		||||
                anchor.x, anchor.y, ksize.width, ksize.height,
 | 
			
		||||
                btype,
 | 
			
		||||
                extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION",
 | 
			
		||||
                isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED");
 | 
			
		||||
 | 
			
		||||
        localsize[0] = BLOCK_SIZE;
 | 
			
		||||
        globalsize[0] = DIVUP(sz.width, BLOCK_SIZE - (ksize.width - 1)) * BLOCK_SIZE;
 | 
			
		||||
        globalsize[1] = DIVUP(sz.height, BLOCK_SIZE_Y);
 | 
			
		||||
 | 
			
		||||
        cv::String errmsg;
 | 
			
		||||
        kernel.create("boxFilter", cv::ocl::imgproc::boxFilter_oclsrc, build_options);
 | 
			
		||||
 | 
			
		||||
        size_t kernelWorkGroupSize = kernel.workGroupSize();
 | 
			
		||||
        if (localsize[0] <= kernelWorkGroupSize)
 | 
			
		||||
            break;
 | 
			
		||||
 | 
			
		||||
        if (BLOCK_SIZE < kernelWorkGroupSize)
 | 
			
		||||
            return false;
 | 
			
		||||
        tryWorkItems = kernelWorkGroupSize;
 | 
			
		||||
    } while (true);
 | 
			
		||||
 | 
			
		||||
    _dst.create(sz, CV_MAKETYPE(ddepth, cn));
 | 
			
		||||
    UMat dst = _dst.getUMat();
 | 
			
		||||
    UMat src = _src.getUMat();
 | 
			
		||||
 | 
			
		||||
    int idxArg = 0;
 | 
			
		||||
    idxArg = kernel.set(idxArg, ocl::KernelArg::ReadOnlyNoSize(src));
 | 
			
		||||
    idxArg = kernel.set(idxArg, ocl::KernelArg::WriteOnly(dst));
 | 
			
		||||
    float borderValue[4] = {0, 0, 0, 0};
 | 
			
		||||
    double borderValueDouble[4] = {0, 0, 0, 0};
 | 
			
		||||
    if ((borderType & ~BORDER_ISOLATED) == BORDER_CONSTANT)
 | 
			
		||||
    {
 | 
			
		||||
        int cnocl = (3 == cn) ? 4 : cn;
 | 
			
		||||
        if (useDouble)
 | 
			
		||||
            idxArg = kernel.set(idxArg, (void *)&borderValueDouble[0], sizeof(double) * cnocl);
 | 
			
		||||
        else
 | 
			
		||||
            idxArg = kernel.set(idxArg, (void *)&borderValue[0], sizeof(float) * cnocl);
 | 
			
		||||
    }
 | 
			
		||||
    if (useDouble)
 | 
			
		||||
        idxArg = kernel.set(idxArg, (double)alpha);
 | 
			
		||||
    else
 | 
			
		||||
        idxArg = kernel.set(idxArg, (float)alpha);
 | 
			
		||||
 | 
			
		||||
    return kernel.run(2, globalsize, localsize, true);
 | 
			
		||||
}
 | 
			
		||||
 | 
			
		||||
}
 | 
			
		||||
 | 
			
		||||
 | 
			
		||||
cv::Ptr<cv::BaseRowFilter> cv::getRowSumFilter(int srcType, int sumType, int ksize, int anchor)
 | 
			
		||||
{
 | 
			
		||||
    int sdepth = CV_MAT_DEPTH(srcType), ddepth = CV_MAT_DEPTH(sumType);
 | 
			
		||||
@@ -712,6 +841,10 @@ void cv::boxFilter( InputArray _src, OutputArray _dst, int ddepth,
 | 
			
		||||
                Size ksize, Point anchor,
 | 
			
		||||
                bool normalize, int borderType )
 | 
			
		||||
{
 | 
			
		||||
    bool use_opencl = ocl::useOpenCL() && _dst.isUMat() && normalize;
 | 
			
		||||
    if( use_opencl && ocl_boxFilter(_src, _dst, ddepth, ksize, anchor, borderType) )
 | 
			
		||||
        return;
 | 
			
		||||
 | 
			
		||||
    Mat src = _src.getMat();
 | 
			
		||||
    int sdepth = src.depth(), cn = src.channels();
 | 
			
		||||
    if( ddepth < 0 )
 | 
			
		||||
 
 | 
			
		||||
							
								
								
									
										132
									
								
								modules/imgproc/test/ocl/test_boxfilter.cpp
									
									
									
									
									
										Normal file
									
								
							
							
						
						
									
										132
									
								
								modules/imgproc/test/ocl/test_boxfilter.cpp
									
									
									
									
									
										Normal file
									
								
							@@ -0,0 +1,132 @@
 | 
			
		||||
/*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-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
 | 
			
		||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
 | 
			
		||||
// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
 | 
			
		||||
// Third party copyrights are property of their respective owners.
 | 
			
		||||
//
 | 
			
		||||
// 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 materials 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/ts/ocl_test.hpp"
 | 
			
		||||
 | 
			
		||||
#ifdef HAVE_OPENCL
 | 
			
		||||
 | 
			
		||||
namespace cvtest {
 | 
			
		||||
namespace ocl {
 | 
			
		||||
 | 
			
		||||
enum
 | 
			
		||||
{
 | 
			
		||||
    noType = -1
 | 
			
		||||
};
 | 
			
		||||
 | 
			
		||||
/////////////////////////////////////////////////////////////////////////////////////////////////
 | 
			
		||||
// boxFilter
 | 
			
		||||
PARAM_TEST_CASE(BoxFilter, MatDepth, Channels, BorderType, bool)
 | 
			
		||||
{
 | 
			
		||||
    int type;
 | 
			
		||||
    Size ksize;
 | 
			
		||||
    Size dsize;
 | 
			
		||||
    Point anchor;
 | 
			
		||||
    int borderType;
 | 
			
		||||
    bool useRoi;
 | 
			
		||||
 | 
			
		||||
    TEST_DECLARE_INPUT_PARAMETER(src)
 | 
			
		||||
    TEST_DECLARE_OUTPUT_PARAMETER(dst)
 | 
			
		||||
 | 
			
		||||
    virtual void SetUp()
 | 
			
		||||
    {
 | 
			
		||||
        type = CV_MAKE_TYPE(GET_PARAM(0), GET_PARAM(1));
 | 
			
		||||
        borderType = GET_PARAM(2);
 | 
			
		||||
        useRoi = GET_PARAM(3);
 | 
			
		||||
    }
 | 
			
		||||
 | 
			
		||||
    void random_roi()
 | 
			
		||||
    {
 | 
			
		||||
        dsize = randomSize(1, MAX_VALUE);
 | 
			
		||||
 | 
			
		||||
        ksize = randomSize(1, dsize.width, 1, dsize.height);
 | 
			
		||||
 | 
			
		||||
        Size roiSize = randomSize(1, MAX_VALUE);
 | 
			
		||||
        Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
 | 
			
		||||
        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
 | 
			
		||||
 | 
			
		||||
        Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
 | 
			
		||||
        randomSubMat(dst, dst_roi, dsize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
 | 
			
		||||
 | 
			
		||||
        anchor.x = anchor.y = -1;
 | 
			
		||||
 | 
			
		||||
        UMAT_UPLOAD_INPUT_PARAMETER(src)
 | 
			
		||||
        UMAT_UPLOAD_OUTPUT_PARAMETER(dst)
 | 
			
		||||
    }
 | 
			
		||||
 | 
			
		||||
    void Near(double threshold = 0.0)
 | 
			
		||||
    {
 | 
			
		||||
        EXPECT_MAT_NEAR(dst, udst, threshold);
 | 
			
		||||
        EXPECT_MAT_NEAR(dst_roi, udst_roi, threshold);
 | 
			
		||||
    }
 | 
			
		||||
};
 | 
			
		||||
 | 
			
		||||
OCL_TEST_P(BoxFilter, Mat)
 | 
			
		||||
{
 | 
			
		||||
    for (int j = 0; j < test_loop_times; j++)
 | 
			
		||||
    {
 | 
			
		||||
        random_roi();
 | 
			
		||||
 | 
			
		||||
        OCL_OFF(cv::boxFilter(src_roi, dst_roi, -1, ksize, anchor, true, borderType));
 | 
			
		||||
        OCL_ON(cv::boxFilter(usrc_roi, udst_roi, -1, ksize, anchor, true, borderType));
 | 
			
		||||
 | 
			
		||||
        Near(1.0);
 | 
			
		||||
    }
 | 
			
		||||
}
 | 
			
		||||
 | 
			
		||||
 | 
			
		||||
OCL_INSTANTIATE_TEST_CASE_P(ImageProc, BoxFilter,
 | 
			
		||||
                            Combine(
 | 
			
		||||
                                Values(CV_8U, CV_16U, CV_16S, CV_32S, CV_32F),
 | 
			
		||||
                                Values(1, 2, 4),
 | 
			
		||||
                                Values((BorderType)BORDER_CONSTANT,
 | 
			
		||||
                                       (BorderType)BORDER_REPLICATE,
 | 
			
		||||
                                       (BorderType)BORDER_REFLECT,
 | 
			
		||||
                                       (BorderType)BORDER_REFLECT_101),
 | 
			
		||||
                                Bool())
 | 
			
		||||
                           );
 | 
			
		||||
 | 
			
		||||
 | 
			
		||||
} } // namespace cvtest::ocl
 | 
			
		||||
 | 
			
		||||
#endif // HAVE_OPENCL
 | 
			
		||||
		Reference in New Issue
	
	Block a user