modules/core/src/opencl/normalize.cl - 4.x - Alx89/OpenCV

modules/core/src/opencl/normalize.cl
72 строки2 KB
Ilya Lavrenov
optimized cv::normalize in case of mask
02 июн 2014, 15:33
02 июн 2014, 15:3317956a5
// This file is part of OpenCV project.
// It is subject to the license terms in the LICENSE file found in the top-level directory
// of this distribution and at http://opencv.org/license.html.

// Copyright (C) 2014, Itseez, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.

#ifdef DOUBLE_SUPPORT
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#elif defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#endif

#define noconvert

#if cn != 3
#define loadpix(addr) *(__global const srcT *)(addr)
#define storepix(val, addr)  *(__global dstT *)(addr) = val
#define srcTSIZE (int)sizeof(srcT)
#define dstTSIZE (int)sizeof(dstT)
#else
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
#define srcTSIZE ((int)sizeof(srcT1)*3)
#define dstTSIZE ((int)sizeof(dstT1)*3)
#endif

__kernel void normalizek(__global const uchar * srcptr, int src_step, int src_offset,
                         __global const uchar * mask, int mask_step, int mask_offset,
                         __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols
#ifdef HAVE_SCALE
                         , float scale
#endif
#ifdef HAVE_DELTA
                         , float delta
#endif
                         )
{
    int x = get_global_id(0);
    int y0 = get_global_id(1) * rowsPerWI;

    if (x < dst_cols)
    {
        int src_index  = mad24(y0, src_step, mad24(x, srcTSIZE, src_offset));
        int mask_index = mad24(y0, mask_step, x + mask_offset);
        int dst_index  = mad24(y0, dst_step, mad24(x, dstTSIZE, dst_offset));

        for (int y = y0, y1 = min(y0 + rowsPerWI, dst_rows); y < y1;
            ++y, src_index += src_step, dst_index += dst_step, mask_index += mask_step)
        {
            if (mask[mask_index])
            {
                workT value = convertToWT(loadpix(srcptr + src_index));
#ifdef HAVE_SCALE
#ifdef HAVE_DELTA
                value = fma(value, (workT)(scale), (workT)(delta));
#else
                value *= (workT)(scale);
#endif
#else // not scale
#ifdef HAVE_DELTA
                value += (workT)(delta);
#endif
#endif

                storepix(convertToDT(value), dstptr + dst_index);
            }
        }
    }
}