← 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/imgproc/src/opencl/bilateral.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/imgproc/src/opencl/bilateral.cl
 
1
//                           License Agreement
 
2
//                For Open Source Computer Vision Library
 
3
//
 
4
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
 
5
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
 
6
// Third party copyrights are property of their respective owners.
 
7
//
 
8
// @Authors
 
9
//    Rock Li, Rock.li@amd.com
 
10
//
 
11
// Redistribution and use in source and binary forms, with or without modification,
 
12
// are permitted provided that the following conditions are met:
 
13
//
 
14
//   * Redistribution's of source code must retain the above copyright notice,
 
15
//     this list of conditions and the following disclaimer.
 
16
//
 
17
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
18
//     this list of conditions and the following disclaimer in the documentation
 
19
//     and/or other materials provided with the distribution.
 
20
//
 
21
//   * The name of the copyright holders may not be used to endorse or promote products
 
22
//     derived from this software without specific prior written permission.
 
23
//
 
24
// This software is provided by the copyright holders and contributors as is and
 
25
// any express or implied warranties, including, but not limited to, the implied
 
26
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
27
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
28
// indirect, incidental, special, exemplary, or consequential damages
 
29
// (including, but not limited to, procurement of substitute goods or services;
 
30
// loss of use, data, or profits; or business interruption) however caused
 
31
// and on any theory of liability, whether in contract, strict liability,
 
32
// or tort (including negligence or otherwise) arising in any way out of
 
33
// the use of this software, even if advised of the possibility of such damage.
 
34
 
 
35
#if cn != 3
 
36
#define loadpix(addr) *(__global const uchar_t *)(addr)
 
37
#define storepix(val, addr)  *(__global uchar_t *)(addr) = val
 
38
#define TSIZE cn
 
39
#else
 
40
#define loadpix(addr) vload3(0, (__global const uchar *)(addr))
 
41
#define storepix(val, addr) vstore3(val, 0, (__global uchar *)(addr))
 
42
#define TSIZE 3
 
43
#endif
 
44
 
 
45
#if cn == 1
 
46
#define SUM(a) a
 
47
#elif cn == 2
 
48
#define SUM(a) a.x + a.y
 
49
#elif cn == 3
 
50
#define SUM(a) a.x + a.y + a.z
 
51
#elif cn == 4
 
52
#define SUM(a) a.x + a.y + a.z + a.w
 
53
#else
 
54
#error "cn should be <= 4"
 
55
#endif
 
56
 
 
57
//Read pixels as integers
 
58
// Intel Device - Read Pixels as floats
 
59
__kernel void bilateral(__global const uchar * src, int src_step, int src_offset,
 
60
                        __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
 
61
                        __constant float * space_weight, __constant int * space_ofs)
 
62
{
 
63
    int x = get_global_id(0);
 
64
    int y = get_global_id(1);
 
65
 
 
66
    if (y < dst_rows && x < dst_cols)
 
67
    {
 
68
        int src_index = mad24(y + radius, src_step, mad24(x + radius, TSIZE, src_offset));
 
69
        int dst_index = mad24(y, dst_step, mad24(x, TSIZE, dst_offset));
 
70
 
 
71
        float_t sum = (float_t)(0.0f);
 
72
        float wsum = 0.0f;
 
73
        #ifdef INTEL_DEVICE
 
74
        float_t val0 = convert_float_t(loadpix(src + src_index));
 
75
        #else
 
76
        int_t val0 = convert_int_t(loadpix(src + src_index));
 
77
        #endif
 
78
        #pragma unroll
 
79
        for (int k = 0; k < maxk; k++ )
 
80
        {
 
81
            #ifdef INTEL_DEVICE
 
82
            float_t val = convert_float_t(loadpix(src + src_index + space_ofs[k]));
 
83
            float diff = SUM(fabs(val - val0));
 
84
            #else
 
85
            int_t val = convert_int_t(loadpix(src + src_index + space_ofs[k]));
 
86
            int diff = SUM(abs(val - val0));
 
87
            #endif
 
88
            float w = space_weight[k] * native_exp((float)(diff * diff * gauss_color_coeff));
 
89
            sum += convert_float_t(val) * (float_t)(w);
 
90
            wsum += w;
 
91
        }
 
92
        storepix(convert_uchar_t(sum / (float_t)(wsum)), dst + dst_index);
 
93
    }
 
94
}
 
95
 
 
96
#ifdef INTEL_DEVICE
 
97
#if cn == 1
 
98
//for single channgel x4 sized images.
 
99
__kernel void bilateral_float4(__global const uchar * src, int src_step, int src_offset,
 
100
                               __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols,
 
101
                               __constant float * space_weight, __constant int * space_ofs)
 
102
{
 
103
    int x = get_global_id(0);
 
104
    int y = get_global_id(1);
 
105
    if (y < dst_rows && x < dst_cols / 4 )
 
106
    {
 
107
        int src_index = ((y + radius) * src_step) + x * 4  + (radius + src_offset);
 
108
        int dst_index = (y  * dst_step) +  x * 4 + dst_offset ;
 
109
        float4 sum = 0.f, wsum = 0.f;
 
110
        float4 val0 = convert_float4(vload4(0, src + src_index));
 
111
        #pragma unroll
 
112
        for (int k = 0; k < maxk; k++ )
 
113
        {
 
114
            float4 val = convert_float4(vload4(0, src + src_index + space_ofs[k]));
 
115
            float4 w = space_weight[k] * native_exp((val - val0) * (val - val0) * gauss_color_coeff);
 
116
            sum += val * w;
 
117
            wsum += w;
 
118
        }
 
119
        sum = sum / wsum + .5f;
 
120
        vstore4(convert_uchar4_rtz(sum), 0, dst + dst_index);
 
121
    }
 
122
}
 
123
#endif
 
124
#endif
 
 
b'\\ No newline at end of file'