← 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/copyset.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/copyset.cl
 
1
/*M///////////////////////////////////////////////////////////////////////////////////////
 
2
//
 
3
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
 
4
//
 
5
//  By downloading, copying, installing or using the software you agree to this license.
 
6
//  If you do not agree to this license, do not download, install,
 
7
//  copy or use the software.
 
8
//
 
9
//
 
10
//                           License Agreement
 
11
//                For Open Source Computer Vision Library
 
12
//
 
13
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
 
14
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
 
15
// Copyright (C) 2013, OpenCV Foundation, all rights reserved.
 
16
// Third party copyrights are property of their respective owners.
 
17
//
 
18
// Redistribution and use in source and binary forms, with or without modification,
 
19
// are permitted provided that the following conditions are met:
 
20
//
 
21
//   * Redistribution's of source code must retain the above copyright notice,
 
22
//     this list of conditions and the following disclaimer.
 
23
//
 
24
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
25
//     this list of conditions and the following disclaimer in the documentation
 
26
//     and/or other materials provided with the distribution.
 
27
//
 
28
//   * The name of the copyright holders may not be used to endorse or promote products
 
29
//     derived from this software without specific prior written permission.
 
30
//
 
31
// This software is provided by the copyright holders and contributors as is and
 
32
// any express or implied warranties, including, but not limited to, the implied
 
33
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
34
// In no event shall the copyright holders or contributors be liable for any direct,
 
35
// indirect, incidental, special, exemplary, or consequential damages
 
36
// (including, but not limited to, procurement of substitute goods or services;
 
37
// loss of use, data, or profits; or business interruption) however caused
 
38
// and on any theory of liability, whether in contract, strict liability,
 
39
// or tort (including negligence or otherwise) arising in any way out of
 
40
// the use of this software, even if advised of the possibility of such damage.
 
41
//
 
42
//M*/
 
43
 
 
44
#ifdef COPY_TO_MASK
 
45
 
 
46
#define DEFINE_DATA \
 
47
    int src_index = mad24(y, src_step, mad24(x, (int)sizeof(T1) * scn, src_offset)); \
 
48
    int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(T1) * scn, dst_offset)); \
 
49
     \
 
50
    __global const T1 * src = (__global const T1 *)(srcptr + src_index); \
 
51
    __global T1 * dst = (__global T1 *)(dstptr + dst_index)
 
52
 
 
53
__kernel void copyToMask(__global const uchar * srcptr, int src_step, int src_offset,
 
54
                         __global const uchar * mask, int mask_step, int mask_offset,
 
55
                         __global uchar * dstptr, int dst_step, int dst_offset,
 
56
                         int dst_rows, int dst_cols)
 
57
{
 
58
    int x = get_global_id(0);
 
59
    int y = get_global_id(1);
 
60
 
 
61
    if (x < dst_cols && y < dst_rows)
 
62
    {
 
63
        mask += mad24(y, mask_step, mad24(x, mcn, mask_offset));
 
64
 
 
65
#if mcn == 1
 
66
        if (mask[0])
 
67
        {
 
68
            DEFINE_DATA;
 
69
 
 
70
            #pragma unroll
 
71
            for (int c = 0; c < scn; ++c)
 
72
                dst[c] = src[c];
 
73
        }
 
74
#ifdef HAVE_DST_UNINIT
 
75
        else
 
76
        {
 
77
            DEFINE_DATA;
 
78
 
 
79
            #pragma unroll
 
80
            for (int c = 0; c < scn; ++c)
 
81
                dst[c] = (T1)(0);
 
82
        }
 
83
#endif
 
84
#elif scn == mcn
 
85
        DEFINE_DATA;
 
86
 
 
87
        #pragma unroll
 
88
        for (int c = 0; c < scn; ++c)
 
89
            if (mask[c])
 
90
                dst[c] = src[c];
 
91
#ifdef HAVE_DST_UNINIT
 
92
            else
 
93
                dst[c] = (T1)(0);
 
94
#endif
 
95
#else
 
96
#error "(mcn == 1 || mcn == scn) should be true"
 
97
#endif
 
98
    }
 
99
}
 
100
 
 
101
#else
 
102
 
 
103
#ifndef dstST
 
104
#define dstST dstT
 
105
#endif
 
106
 
 
107
#if cn != 3
 
108
#define value value_
 
109
#define storedst(val) *(__global dstT *)(dstptr + dst_index) = val
 
110
#else
 
111
#define value (dstT)(value_.x, value_.y, value_.z)
 
112
#define storedst(val) vstore3(val, 0, (__global dstT1 *)(dstptr + dst_index))
 
113
#endif
 
114
 
 
115
__kernel void setMask(__global const uchar* mask, int maskstep, int maskoffset,
 
116
                      __global uchar* dstptr, int dststep, int dstoffset,
 
117
                      int rows, int cols, dstST value_)
 
118
{
 
119
    int x = get_global_id(0);
 
120
    int y0 = get_global_id(1) * rowsPerWI;
 
121
 
 
122
    if (x < cols)
 
123
    {
 
124
        int mask_index = mad24(y0, maskstep, x + maskoffset);
 
125
        int dst_index  = mad24(y0, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset));
 
126
 
 
127
        for (int y = y0, y1 = min(rows, y0 + rowsPerWI); y < y1; ++y)
 
128
        {
 
129
            if( mask[mask_index] )
 
130
                storedst(value);
 
131
 
 
132
            mask_index += maskstep;
 
133
            dst_index += dststep;
 
134
        }
 
135
    }
 
136
}
 
137
 
 
138
__kernel void set(__global uchar* dstptr, int dststep, int dstoffset,
 
139
                  int rows, int cols, dstST value_)
 
140
{
 
141
    int x = get_global_id(0);
 
142
    int y0 = get_global_id(1) * rowsPerWI;
 
143
 
 
144
    if (x < cols)
 
145
    {
 
146
        int dst_index  = mad24(y0, dststep, mad24(x, (int)sizeof(dstT1) * cn, dstoffset));
 
147
 
 
148
        for (int y = y0, y1 = min(rows, y0 + rowsPerWI); y < y1; ++y, dst_index += dststep)
 
149
            storedst(value);
 
150
    }
 
151
}
 
152
 
 
153
#endif