← Back to branch summary

~paparazzi-uav/paparazzi/v5.0-manual

« back to all changes in this revision

Viewing changes to sw/ext/opencv_bebop/opencv/modules/core/src/opencl/normalize.cl

  • Committer: Paparazzi buildbot
  • Date: 2016-05-18 15:00:29 UTC
  • Revision ID: felix.ruess+docbot@gmail.com-20160518150029-e8lgzi5kvb4p7un9
Manual import commit 4b8bbb730080dac23cf816b98908dacfabe2a8ec from v5.0 branch.

Show diffs side-by-side

added added

removed removed

Lines of Context:
sw/ext/opencv_bebop/opencv/modules/core/src/opencl/normalize.cl
 
1
// This file is part of OpenCV project.
 
2
// It is subject to the license terms in the LICENSE file found in the top-level directory
 
3
// of this distribution and at http://opencv.org/license.html.
 
4
 
 
5
// Copyright (C) 2014, Itseez, Inc., all rights reserved.
 
6
// Third party copyrights are property of their respective owners.
 
7
 
 
8
#ifdef DOUBLE_SUPPORT
 
9
#ifdef cl_amd_fp64
 
10
#pragma OPENCL EXTENSION cl_amd_fp64:enable
 
11
#elif defined (cl_khr_fp64)
 
12
#pragma OPENCL EXTENSION cl_khr_fp64:enable
 
13
#endif
 
14
#endif
 
15
 
 
16
#define noconvert
 
17
 
 
18
#if cn != 3
 
19
#define loadpix(addr) *(__global const srcT *)(addr)
 
20
#define storepix(val, addr)  *(__global dstT *)(addr) = val
 
21
#define srcTSIZE (int)sizeof(srcT)
 
22
#define dstTSIZE (int)sizeof(dstT)
 
23
#else
 
24
#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr))
 
25
#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr))
 
26
#define srcTSIZE ((int)sizeof(srcT1)*3)
 
27
#define dstTSIZE ((int)sizeof(dstT1)*3)
 
28
#endif
 
29
 
 
30
__kernel void normalizek(__global const uchar * srcptr, int src_step, int src_offset,
 
31
                         __global const uchar * mask, int mask_step, int mask_offset,
 
32
                         __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols
 
33
#ifdef HAVE_SCALE
 
34
                         , float scale
 
35
#endif
 
36
#ifdef HAVE_DELTA
 
37
                         , float delta
 
38
#endif
 
39
                         )
 
40
{
 
41
    int x = get_global_id(0);
 
42
    int y0 = get_global_id(1) * rowsPerWI;
 
43
 
 
44
    if (x < dst_cols)
 
45
    {
 
46
        int src_index  = mad24(y0, src_step, mad24(x, srcTSIZE, src_offset));
 
47
        int mask_index = mad24(y0, mask_step, x + mask_offset);
 
48
        int dst_index  = mad24(y0, dst_step, mad24(x, dstTSIZE, dst_offset));
 
49
 
 
50
        for (int y = y0, y1 = min(y0 + rowsPerWI, dst_rows); y < y1;
 
51
            ++y, src_index += src_step, dst_index += dst_step, mask_index += mask_step)
 
52
        {
 
53
            if (mask[mask_index])
 
54
            {
 
55
                workT value = convertToWT(loadpix(srcptr + src_index));
 
56
#ifdef HAVE_SCALE
 
57
#ifdef HAVE_DELTA
 
58
                value = fma(value, (workT)(scale), (workT)(delta));
 
59
#else
 
60
                value *= (workT)(scale);
 
61
#endif
 
62
#else // not scale
 
63
#ifdef HAVE_DELTA
 
64
                value += (workT)(delta);
 
65
#endif
 
66
#endif
 
67
 
 
68
                storepix(convertToDT(value), dstptr + dst_index);
 
69
            }
 
70
        }
 
71
    }
 
72
}