CoCalc -- multibandblend.cl

GitHub Repository: Tetragramm/opencv
Path: blob/master/modules/stitching/src/opencl/multibandblend.cl
¹⁶³³⁷ views
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

7
//
8
// Common preprocessors macro
9
//
10

11
//
12
// TODO: Move this common code into "header" file
13
//
14

15
#ifndef NL // New Line: for preprocessor debugging
16
#define NL
17
#endif
18

19
#define REF(x) x
20
#define __CAT(x, y) x##y
21
#define CAT(x, y) __CAT(x, y)
22

23
//
24
// All matrixes are come with this description ("name" is a name of matrix):
25
// * name_CN - number of channels (1,2,3,4)
26
// * name_DEPTH - numeric value of CV_MAT_DEPTH(type). See CV_8U, CV_32S, etc macro below.
27
//
28
// Currently we also pass these attributes (to reduce this macro block):
29
// * name_T - datatype (int, float, uchar4, float4)
30
// * name_T1 - datatype for one channel (int, float, uchar).
31
//   It is equal to result of "T1(name_T)" macro
32
// * name_TSIZE - CV_ELEM_SIZE(type).
33
//   We can't use sizeof(name_T) here, because sizeof(float3) is usually equal to 8, not 6.
34
// * name_T1SIZE - CV_ELEM_SIZE1(type)
35
//
36

37
//
38
// Usage sample:
39
//
40
// #define workType TYPE(float, src_CN)
41
// #define convertToWorkType CONVERT_TO(workType)
42
// #define convertWorkTypeToDstType CONVERT(workType, dst_T)
43
//
44
// __kernel void kernelFn(DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(dst))
45
// {
46
//     const int x = get_global_id(0);
47
//     const int y = get_global_id(1);
48
//
49
//     if (x < srcWidth && y < srcHeight)
50
//     {
51
//         int src_byteOffset = MAT_BYTE_OFFSET(src, x, y);
52
//         int dst_byteOffset = MAT_BYTE_OFFSET(dst, x, y);
53
//         workType value = convertToWorkType(LOAD_MAT_AT(src, src_byteOffset));
54
//
55
//         ... value processing ...
56
//
57
//         STORE_MAT_AT(dst, dst_byteOffset, convertWorkTypeToDstType(value));
58
//     }
59
// }
60
//
61

62
#define DECLARE_MAT_ARG(name) \
63
    __global uchar* restrict name ## Ptr, \
64
    int name ## StepBytes, \
65
    int name ## Offset, \
66
    int name ## Height, \
67
    int name ## Width NL
68

69
#define MAT_BYTE_OFFSET(name, x, y) mad24((y)/* + name ## OffsetY*/, name ## StepBytes, ((x)/* + name ## OffsetX*/) * (int)(name ## _TSIZE) + name ## Offset)
70
#define MAT_RELATIVE_BYTE_OFFSET(name, x, y) mad24(y, name ## StepBytes, (x) * (int)(name ## _TSIZE))
71

72
#define __LOAD_MAT_AT(name, byteOffset) *((const __global name ## _T*)(name ## Ptr + (byteOffset)))
73
#define __vload_CN__(name_cn) vload ## name_cn
74
#define __vload_CN_(name_cn) __vload_CN__(name_cn)
75
#define __vload_CN(name) __vload_CN_(name ## _CN)
76
#define __LOAD_MAT_AT_vload(name, byteOffset) __vload_CN(name)(0, ((const __global name ## _T1*)(name ## Ptr + (byteOffset))))
77
#define __LOAD_MAT_AT_1 __LOAD_MAT_AT
78
#define __LOAD_MAT_AT_2 __LOAD_MAT_AT
79
#define __LOAD_MAT_AT_3 __LOAD_MAT_AT_vload
80
#define __LOAD_MAT_AT_4 __LOAD_MAT_AT
81
#define __LOAD_MAT_AT_CN__(name_cn) __LOAD_MAT_AT_ ## name_cn
82
#define __LOAD_MAT_AT_CN_(name_cn) __LOAD_MAT_AT_CN__(name_cn)
83
#define __LOAD_MAT_AT_CN(name) __LOAD_MAT_AT_CN_(name ## _CN)
84
#define LOAD_MAT_AT(name, byteOffset) __LOAD_MAT_AT_CN(name)(name, byteOffset)
85

86
#define __STORE_MAT_AT(name, byteOffset, v) *((__global name ## _T*)(name ## Ptr + (byteOffset))) = v
87
#define __vstore_CN__(name_cn) vstore ## name_cn
88
#define __vstore_CN_(name_cn) __vstore_CN__(name_cn)
89
#define __vstore_CN(name) __vstore_CN_(name ## _CN)
90
#define __STORE_MAT_AT_vstore(name, byteOffset, v) __vstore_CN(name)(v, 0, ((__global name ## _T1*)(name ## Ptr + (byteOffset))))
91
#define __STORE_MAT_AT_1 __STORE_MAT_AT
92
#define __STORE_MAT_AT_2 __STORE_MAT_AT
93
#define __STORE_MAT_AT_3 __STORE_MAT_AT_vstore
94
#define __STORE_MAT_AT_4 __STORE_MAT_AT
95
#define __STORE_MAT_AT_CN__(name_cn) __STORE_MAT_AT_ ## name_cn
96
#define __STORE_MAT_AT_CN_(name_cn) __STORE_MAT_AT_CN__(name_cn)
97
#define __STORE_MAT_AT_CN(name) __STORE_MAT_AT_CN_(name ## _CN)
98
#define STORE_MAT_AT(name, byteOffset, v) __STORE_MAT_AT_CN(name)(name, byteOffset, v)
99

100
#define T1_uchar uchar
101
#define T1_uchar2 uchar
102
#define T1_uchar3 uchar
103
#define T1_uchar4 uchar
104
#define T1_char char
105
#define T1_char2 char
106
#define T1_char3 char
107
#define T1_char4 char
108
#define T1_ushort ushort
109
#define T1_ushort2 ushort
110
#define T1_ushort3 ushort
111
#define T1_ushort4 ushort
112
#define T1_short short
113
#define T1_short2 short
114
#define T1_short3 short
115
#define T1_short4 short
116
#define T1_int int
117
#define T1_int2 int
118
#define T1_int3 int
119
#define T1_int4 int
120
#define T1_float float
121
#define T1_float2 float
122
#define T1_float3 float
123
#define T1_float4 float
124
#define T1_double double
125
#define T1_double2 double
126
#define T1_double3 double
127
#define T1_double4 double
128
#define T1(type) REF(CAT(T1_, REF(type)))
129

130
#define uchar1 uchar
131
#define char1 char
132
#define short1 short
133
#define ushort1 ushort
134
#define int1 int
135
#define float1 float
136
#define double1 double
137
#define TYPE(type, cn) REF(CAT(REF(type), REF(cn)))
138

139
#define __CONVERT_MODE_uchar_uchar __NO_CONVERT
140
#define __CONVERT_MODE_uchar_char __CONVERT_sat
141
#define __CONVERT_MODE_uchar_ushort __CONVERT
142
#define __CONVERT_MODE_uchar_short __CONVERT
143
#define __CONVERT_MODE_uchar_int __CONVERT
144
#define __CONVERT_MODE_uchar_float __CONVERT
145
#define __CONVERT_MODE_uchar_double __CONVERT
146
#define __CONVERT_MODE_char_uchar __CONVERT_sat
147
#define __CONVERT_MODE_char_char __NO_CONVERT
148
#define __CONVERT_MODE_char_ushort __CONVERT_sat
149
#define __CONVERT_MODE_char_short __CONVERT
150
#define __CONVERT_MODE_char_int __CONVERT
151
#define __CONVERT_MODE_char_float __CONVERT
152
#define __CONVERT_MODE_char_double __CONVERT
153
#define __CONVERT_MODE_ushort_uchar __CONVERT_sat
154
#define __CONVERT_MODE_ushort_char __CONVERT_sat
155
#define __CONVERT_MODE_ushort_ushort __NO_CONVERT
156
#define __CONVERT_MODE_ushort_short __CONVERT_sat
157
#define __CONVERT_MODE_ushort_int __CONVERT
158
#define __CONVERT_MODE_ushort_float __CONVERT
159
#define __CONVERT_MODE_ushort_double __CONVERT
160
#define __CONVERT_MODE_short_uchar __CONVERT_sat
161
#define __CONVERT_MODE_short_char __CONVERT_sat
162
#define __CONVERT_MODE_short_ushort __CONVERT_sat
163
#define __CONVERT_MODE_short_short __NO_CONVERT
164
#define __CONVERT_MODE_short_int __CONVERT
165
#define __CONVERT_MODE_short_float __CONVERT
166
#define __CONVERT_MODE_short_double __CONVERT
167
#define __CONVERT_MODE_int_uchar __CONVERT_sat
168
#define __CONVERT_MODE_int_char __CONVERT_sat
169
#define __CONVERT_MODE_int_ushort __CONVERT_sat
170
#define __CONVERT_MODE_int_short __CONVERT_sat
171
#define __CONVERT_MODE_int_int __NO_CONVERT
172
#define __CONVERT_MODE_int_float __CONVERT
173
#define __CONVERT_MODE_int_double __CONVERT
174
#define __CONVERT_MODE_float_uchar __CONVERT_sat_rte
175
#define __CONVERT_MODE_float_char __CONVERT_sat_rte
176
#define __CONVERT_MODE_float_ushort __CONVERT_sat_rte
177
#define __CONVERT_MODE_float_short __CONVERT_sat_rte
178
#define __CONVERT_MODE_float_int __CONVERT_rte
179
#define __CONVERT_MODE_float_float __NO_CONVERT
180
#define __CONVERT_MODE_float_double __CONVERT
181
#define __CONVERT_MODE_double_uchar __CONVERT_sat_rte
182
#define __CONVERT_MODE_double_char __CONVERT_sat_rte
183
#define __CONVERT_MODE_double_ushort __CONVERT_sat_rte
184
#define __CONVERT_MODE_double_short __CONVERT_sat_rte
185
#define __CONVERT_MODE_double_int __CONVERT_rte
186
#define __CONVERT_MODE_double_float __CONVERT
187
#define __CONVERT_MODE_double_double __NO_CONVERT
188
#define __CONVERT_MODE(srcType, dstType) CAT(__CONVERT_MODE_, CAT(REF(T1(srcType)), CAT(_, REF(T1(dstType)))))
189

190
#define __ROUND_MODE__NO_CONVERT
191
#define __ROUND_MODE__CONVERT // nothing
192
#define __ROUND_MODE__CONVERT_rte _rte
193
#define __ROUND_MODE__CONVERT_sat _sat
194
#define __ROUND_MODE__CONVERT_sat_rte _sat_rte
195
#define ROUND_MODE(srcType, dstType) CAT(__ROUND_MODE_, __CONVERT_MODE(srcType, dstType))
196

197
#define __CONVERT_ROUND(dstType, roundMode) CAT(CAT(convert_, REF(dstType)), roundMode)
198
#define __NO_CONVERT(dstType) // nothing
199
#define __CONVERT(dstType) __CONVERT_ROUND(dstType,)
200
#define __CONVERT_rte(dstType) __CONVERT_ROUND(dstType,_rte)
201
#define __CONVERT_sat(dstType) __CONVERT_ROUND(dstType,_sat)
202
#define __CONVERT_sat_rte(dstType) __CONVERT_ROUND(dstType,_sat_rte)
203
#define CONVERT(srcType, dstType) REF(__CONVERT_MODE(srcType,dstType))(dstType)
204
#define CONVERT_TO(dstType) __CONVERT_ROUND(dstType,)
205

206
// OpenCV depths
207
#define CV_8U   0
208
#define CV_8S   1
209
#define CV_16U  2
210
#define CV_16S  3
211
#define CV_32S  4
212
#define CV_32F  5
213
#define CV_64F  6
214

215
//
216
// End of common preprocessors macro
217
//
218

219

220

221
#if defined(DEFINE_feed)
222

223
#define workType TYPE(weight_T1, src_CN)
224

225
#if src_DEPTH == 3 && src_CN == 3
226
#define convertSrcToWorkType convert_float3
227
#else
228
#define convertSrcToWorkType CONVERT_TO(workType)
229
#endif
230

231
#if dst_DEPTH == 3 && dst_CN == 3
232
#define convertToDstType convert_short3
233
#else
234
#define convertToDstType CONVERT_TO(dst_T) // sat_rte provides incompatible results with CPU path
235
#endif
236

237
__kernel void feed(
238
        DECLARE_MAT_ARG(src), DECLARE_MAT_ARG(weight),
239
        DECLARE_MAT_ARG(dst), DECLARE_MAT_ARG(dstWeight)
240
)
241
{
242
    const int x = get_global_id(0);
243
    const int y = get_global_id(1);
244

245
    if (x < srcWidth && y < srcHeight)
246
    {
247
        int src_byteOffset = MAT_BYTE_OFFSET(src, x, y);
248
        int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y);
249
        int dst_byteOffset = MAT_BYTE_OFFSET(dst, x, y);
250
        int dstWeight_byteOffset = MAT_BYTE_OFFSET(dstWeight, x, y);
251

252
        weight_T w = LOAD_MAT_AT(weight, weight_byteOffset);
253
        workType src_value = convertSrcToWorkType(LOAD_MAT_AT(src, src_byteOffset));
254
        STORE_MAT_AT(dst, dst_byteOffset, LOAD_MAT_AT(dst, dst_byteOffset) + convertToDstType(src_value * w));
255
        STORE_MAT_AT(dstWeight, dstWeight_byteOffset, LOAD_MAT_AT(dstWeight, dstWeight_byteOffset) + w);
256
    }
257
}
258

259
#endif
260

261
#if defined(DEFINE_normalizeUsingWeightMap)
262

263
#if mat_DEPTH == 3 && mat_CN == 3
264
#define workType float3
265
#define convertSrcToWorkType convert_float3
266
#define convertToDstType convert_short3
267
#else
268
#define workType TYPE(weight_T1, mat_CN)
269
#define convertSrcToWorkType CONVERT_TO(workType)
270
#define convertToDstType CONVERT_TO(mat_T) // sat_rte provides incompatible results with CPU path
271
#endif
272

273
#if weight_DEPTH >= CV_32F
274
#define WEIGHT_EPS 1e-5f
275
#else
276
#define WEIGHT_EPS 0
277
#endif
278

279
__kernel void normalizeUsingWeightMap(
280
        DECLARE_MAT_ARG(mat), DECLARE_MAT_ARG(weight)
281
)
282
{
283
    const int x = get_global_id(0);
284
    const int y = get_global_id(1);
285

286
    if (x < matWidth && y < matHeight)
287
    {
288
        int mat_byteOffset = MAT_BYTE_OFFSET(mat, x, y);
289
        int weight_byteOffset = MAT_BYTE_OFFSET(weight, x, y);
290

291
        weight_T w = LOAD_MAT_AT(weight, weight_byteOffset);
292
        workType value = convertSrcToWorkType(LOAD_MAT_AT(mat, mat_byteOffset));
293
        value = value / (w + WEIGHT_EPS);
294
        STORE_MAT_AT(mat, mat_byteOffset, convertToDstType(value));
295
    }
296
}
297

298
#endif
299

300
Product

Resources

Company
