← 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/cvtcolor.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/cvtcolor.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
// Third party copyrights are property of their respective owners.
 
16
//
 
17
// @Authors
 
18
//    Jia Haipeng, jiahaipeng95@gmail.com
 
19
//    Peng Xiao, pengxiao@multicorewareinc.com
 
20
//
 
21
// Redistribution and use in source and binary forms, with or without modification,
 
22
// are permitted provided that the following conditions are met:
 
23
//
 
24
//   * Redistribution's of source code must retain the above copyright notice,
 
25
//     this list of conditions and the following disclaimer.
 
26
//
 
27
//   * Redistribution's in binary form must reproduce the above copyright notice,
 
28
//     this list of conditions and the following disclaimer in the documentation
 
29
//     and/or other materials provided with the distribution.
 
30
//
 
31
//   * The name of the copyright holders may not be used to endorse or promote products
 
32
//     derived from this software without specific prior written permission.
 
33
//
 
34
// This software is provided by the copyright holders and contributors as is and
 
35
// any express or implied warranties, including, but not limited to, the implied
 
36
// warranties of merchantability and fitness for a particular purpose are disclaimed.
 
37
// In no event shall the Intel Corporation or contributors be liable for any direct,
 
38
// indirect, incidental, special, exemplary, or consequential damages
 
39
// (including, but not limited to, procurement of substitute goods or services;
 
40
// loss of use, data, or profits; or business interruption) however caused
 
41
// and on any theory of liability, whether in contract, strict liability,
 
42
// or tort (including negligence or otherwise) arising in any way out of
 
43
// the use of this software, even if advised of the possibility of such damage.
 
44
//
 
45
//M*/
 
46
 
 
47
/**************************************PUBLICFUNC*************************************/
 
48
 
 
49
#if depth == 0
 
50
    #define DATA_TYPE uchar
 
51
    #define MAX_NUM  255
 
52
    #define HALF_MAX 128
 
53
    #define COEFF_TYPE int
 
54
    #define SAT_CAST(num) convert_uchar_sat(num)
 
55
    #define DEPTH_0
 
56
#elif depth == 2
 
57
    #define DATA_TYPE ushort
 
58
    #define MAX_NUM  65535
 
59
    #define HALF_MAX 32768
 
60
    #define COEFF_TYPE int
 
61
    #define SAT_CAST(num) convert_ushort_sat(num)
 
62
    #define DEPTH_2
 
63
#elif depth == 5
 
64
    #define DATA_TYPE float
 
65
    #define MAX_NUM  1.0f
 
66
    #define HALF_MAX 0.5f
 
67
    #define COEFF_TYPE float
 
68
    #define SAT_CAST(num) (num)
 
69
    #define DEPTH_5
 
70
#else
 
71
    #error "invalid depth: should be 0 (CV_8U), 2 (CV_16U) or 5 (CV_32F)"
 
72
#endif
 
73
 
 
74
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
 
75
 
 
76
enum
 
77
{
 
78
    yuv_shift  = 14,
 
79
    xyz_shift  = 12,
 
80
    hsv_shift  = 12,
 
81
    R2Y        = 4899,
 
82
    G2Y        = 9617,
 
83
    B2Y        = 1868,
 
84
    BLOCK_SIZE = 256
 
85
};
 
86
 
 
87
#define scnbytes ((int)sizeof(DATA_TYPE)*scn)
 
88
#define dcnbytes ((int)sizeof(DATA_TYPE)*dcn)
 
89
 
 
90
#ifndef hscale
 
91
#define hscale 0
 
92
#endif
 
93
 
 
94
#ifndef hrange
 
95
#define hrange 0
 
96
#endif
 
97
 
 
98
#if bidx == 0
 
99
#define R_COMP z
 
100
#define G_COMP y
 
101
#define B_COMP x
 
102
#elif bidx == 2
 
103
#define R_COMP x
 
104
#define G_COMP y
 
105
#define B_COMP z
 
106
#elif bidx == 3
 
107
// The only kernel that uses bidx == 3 doesn't use these macros.
 
108
// But we still need to make the compiler happy.
 
109
#define R_COMP w
 
110
#define G_COMP w
 
111
#define B_COMP w
 
112
#endif
 
113
 
 
114
#ifndef uidx
 
115
#define uidx 0
 
116
#endif
 
117
 
 
118
#ifndef yidx
 
119
#define yidx 0
 
120
#endif
 
121
 
 
122
#ifndef PIX_PER_WI_X
 
123
#define PIX_PER_WI_X 1
 
124
#endif
 
125
 
 
126
#define __CAT(x, y) x##y
 
127
#define CAT(x, y) __CAT(x, y)
 
128
 
 
129
#define DATA_TYPE_4 CAT(DATA_TYPE, 4)
 
130
 
 
131
///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
 
132
 
 
133
__kernel void RGB2Gray(__global const uchar * srcptr, int src_step, int src_offset,
 
134
                       __global uchar * dstptr, int dst_step, int dst_offset,
 
135
                       int rows, int cols)
 
136
{
 
137
    int x = get_global_id(0);
 
138
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
139
 
 
140
    if (x < cols)
 
141
    {
 
142
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
143
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
144
 
 
145
        #pragma unroll
 
146
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
147
        {
 
148
            if (y < rows)
 
149
            {
 
150
                __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
 
151
                __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
 
152
                DATA_TYPE_4 src_pix = vload4(0, src);
 
153
#ifdef DEPTH_5
 
154
                dst[0] = fma(src_pix.B_COMP, 0.114f, fma(src_pix.G_COMP, 0.587f, src_pix.R_COMP * 0.299f));
 
155
#else
 
156
                dst[0] = (DATA_TYPE)CV_DESCALE(mad24(src_pix.B_COMP, B2Y, mad24(src_pix.G_COMP, G2Y, mul24(src_pix.R_COMP, R2Y))), yuv_shift);
 
157
#endif
 
158
                ++y;
 
159
                src_index += src_step;
 
160
                dst_index += dst_step;
 
161
            }
 
162
        }
 
163
    }
 
164
}
 
165
 
 
166
__kernel void Gray2RGB(__global const uchar * srcptr, int src_step, int src_offset,
 
167
                       __global uchar * dstptr, int dst_step, int dst_offset,
 
168
                       int rows, int cols)
 
169
{
 
170
    int x = get_global_id(0);
 
171
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
172
 
 
173
    if (x < cols)
 
174
    {
 
175
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
176
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
177
 
 
178
        #pragma unroll
 
179
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
180
        {
 
181
            if (y < rows)
 
182
            {
 
183
                __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
 
184
                __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
 
185
                DATA_TYPE val = src[0];
 
186
#if dcn == 3 || defined DEPTH_5
 
187
                dst[0] = dst[1] = dst[2] = val;
 
188
#if dcn == 4
 
189
                dst[3] = MAX_NUM;
 
190
#endif
 
191
#else
 
192
                *(__global DATA_TYPE_4 *)dst = (DATA_TYPE_4)(val, val, val, MAX_NUM);
 
193
#endif
 
194
                ++y;
 
195
                dst_index += dst_step;
 
196
                src_index += src_step;
 
197
            }
 
198
        }
 
199
    }
 
200
}
 
201
 
 
202
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
 
203
 
 
204
__constant float c_RGB2YUVCoeffs_f[5]  = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
 
205
__constant int   c_RGB2YUVCoeffs_i[5]  = { B2Y, G2Y, R2Y, 8061, 14369 };
 
206
 
 
207
__kernel void RGB2YUV(__global const uchar* srcptr, int src_step, int src_offset,
 
208
                      __global uchar* dstptr, int dst_step, int dt_offset,
 
209
                      int rows, int cols)
 
210
{
 
211
    int x = get_global_id(0);
 
212
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
213
 
 
214
    if (x < cols)
 
215
    {
 
216
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
217
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset));
 
218
 
 
219
        #pragma unroll
 
220
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
221
        {
 
222
            if (y < rows)
 
223
            {
 
224
                __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
 
225
                __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
 
226
                DATA_TYPE_4 src_pix = vload4(0, src);
 
227
                DATA_TYPE b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
228
 
 
229
#ifdef DEPTH_5
 
230
                __constant float * coeffs = c_RGB2YUVCoeffs_f;
 
231
                const DATA_TYPE Y = fma(b, coeffs[0], fma(g, coeffs[1], r * coeffs[2]));
 
232
                const DATA_TYPE U = fma(b - Y, coeffs[3], HALF_MAX);
 
233
                const DATA_TYPE V = fma(r - Y, coeffs[4], HALF_MAX);
 
234
#else
 
235
                __constant int * coeffs = c_RGB2YUVCoeffs_i;
 
236
                const int delta = HALF_MAX * (1 << yuv_shift);
 
237
                const int Y = CV_DESCALE(mad24(b, coeffs[0], mad24(g, coeffs[1], mul24(r, coeffs[2]))), yuv_shift);
 
238
                const int U = CV_DESCALE(mad24(b - Y, coeffs[3], delta), yuv_shift);
 
239
                const int V = CV_DESCALE(mad24(r - Y, coeffs[4], delta), yuv_shift);
 
240
#endif
 
241
 
 
242
                dst[0] = SAT_CAST( Y );
 
243
                dst[1] = SAT_CAST( U );
 
244
                dst[2] = SAT_CAST( V );
 
245
 
 
246
                ++y;
 
247
                dst_index += dst_step;
 
248
                src_index += src_step;
 
249
            }
 
250
        }
 
251
    }
 
252
}
 
253
 
 
254
__constant float c_YUV2RGBCoeffs_f[4] = { 2.032f, -0.395f, -0.581f, 1.140f };
 
255
__constant int   c_YUV2RGBCoeffs_i[4] = { 33292, -6472, -9519, 18678 };
 
256
 
 
257
__kernel void YUV2RGB(__global const uchar* srcptr, int src_step, int src_offset,
 
258
                      __global uchar* dstptr, int dst_step, int dt_offset,
 
259
                      int rows, int cols)
 
260
{
 
261
    int x = get_global_id(0);
 
262
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
263
 
 
264
    if (x < cols)
 
265
    {
 
266
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
267
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset));
 
268
 
 
269
        #pragma unroll
 
270
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
271
        {
 
272
            if (y < rows)
 
273
            {
 
274
                __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
 
275
                __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
 
276
                DATA_TYPE_4 src_pix = vload4(0, src);
 
277
                DATA_TYPE Y = src_pix.x, U = src_pix.y, V = src_pix.z;
 
278
 
 
279
#ifdef DEPTH_5
 
280
                __constant float * coeffs = c_YUV2RGBCoeffs_f;
 
281
                float r = fma(V - HALF_MAX, coeffs[3], Y);
 
282
                float g = fma(V - HALF_MAX, coeffs[2], fma(U - HALF_MAX, coeffs[1], Y));
 
283
                float b = fma(U - HALF_MAX, coeffs[0], Y);
 
284
#else
 
285
                __constant int * coeffs = c_YUV2RGBCoeffs_i;
 
286
                const int r = Y + CV_DESCALE(mul24(V - HALF_MAX, coeffs[3]), yuv_shift);
 
287
                const int g = Y + CV_DESCALE(mad24(V - HALF_MAX, coeffs[2], mul24(U - HALF_MAX, coeffs[1])), yuv_shift);
 
288
                const int b = Y + CV_DESCALE(mul24(U - HALF_MAX, coeffs[0]), yuv_shift);
 
289
#endif
 
290
 
 
291
                dst[bidx] = SAT_CAST( b );
 
292
                dst[1] = SAT_CAST( g );
 
293
                dst[bidx^2] = SAT_CAST( r );
 
294
#if dcn == 4
 
295
                dst[3] = MAX_NUM;
 
296
#endif
 
297
                ++y;
 
298
                dst_index += dst_step;
 
299
                src_index += src_step;
 
300
            }
 
301
        }
 
302
    }
 
303
}
 
304
__constant float c_YUV2RGBCoeffs_420[5] = { 1.163999557f, 2.017999649f, -0.390999794f,
 
305
                                            -0.812999725f, 1.5959997177f };
 
306
 
 
307
__kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset,
 
308
                            __global uchar* dstptr, int dst_step, int dt_offset,
 
309
                            int rows, int cols)
 
310
{
 
311
    int x = get_global_id(0);
 
312
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
313
 
 
314
    if (x < cols / 2)
 
315
    {
 
316
        #pragma unroll
 
317
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
318
        {
 
319
            if (y < rows / 2 )
 
320
            {
 
321
                __global const uchar* ysrc = srcptr + mad24(y << 1, src_step, (x << 1) + src_offset);
 
322
                __global const uchar* usrc = srcptr + mad24(rows + y, src_step, (x << 1) + src_offset);
 
323
                __global uchar*       dst1 = dstptr + mad24(y << 1, dst_step, mad24(x, dcn<<1, dt_offset));
 
324
                __global uchar*       dst2 = dst1 + dst_step;
 
325
 
 
326
                float Y1 = ysrc[0];
 
327
                float Y2 = ysrc[1];
 
328
                float Y3 = ysrc[src_step];
 
329
                float Y4 = ysrc[src_step + 1];
 
330
 
 
331
                float U  = ((float)usrc[uidx]) - HALF_MAX;
 
332
                float V  = ((float)usrc[1-uidx]) - HALF_MAX;
 
333
 
 
334
                __constant float* coeffs = c_YUV2RGBCoeffs_420;
 
335
                float ruv = fma(coeffs[4], V, 0.5f);
 
336
                float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f));
 
337
                float buv = fma(coeffs[1], U, 0.5f);
 
338
 
 
339
                Y1 = max(0.f, Y1 - 16.f) * coeffs[0];
 
340
                dst1[2 - bidx] = convert_uchar_sat(Y1 + ruv);
 
341
                dst1[1]        = convert_uchar_sat(Y1 + guv);
 
342
                dst1[bidx]     = convert_uchar_sat(Y1 + buv);
 
343
#if dcn == 4
 
344
                dst1[3]        = 255;
 
345
#endif
 
346
 
 
347
                Y2 = max(0.f, Y2 - 16.f) * coeffs[0];
 
348
                dst1[dcn + 2 - bidx] = convert_uchar_sat(Y2 + ruv);
 
349
                dst1[dcn + 1]        = convert_uchar_sat(Y2 + guv);
 
350
                dst1[dcn + bidx]     = convert_uchar_sat(Y2 + buv);
 
351
#if dcn == 4
 
352
                dst1[7]        = 255;
 
353
#endif
 
354
 
 
355
                Y3 = max(0.f, Y3 - 16.f) * coeffs[0];
 
356
                dst2[2 - bidx] = convert_uchar_sat(Y3 + ruv);
 
357
                dst2[1]        = convert_uchar_sat(Y3 + guv);
 
358
                dst2[bidx]     = convert_uchar_sat(Y3 + buv);
 
359
#if dcn == 4
 
360
                dst2[3]        = 255;
 
361
#endif
 
362
 
 
363
                Y4 = max(0.f, Y4 - 16.f) * coeffs[0];
 
364
                dst2[dcn + 2 - bidx] = convert_uchar_sat(Y4 + ruv);
 
365
                dst2[dcn + 1]        = convert_uchar_sat(Y4 + guv);
 
366
                dst2[dcn + bidx]     = convert_uchar_sat(Y4 + buv);
 
367
#if dcn == 4
 
368
                dst2[7]        = 255;
 
369
#endif
 
370
            }
 
371
            ++y;
 
372
        }
 
373
    }
 
374
}
 
375
 
 
376
__kernel void YUV2RGB_YV12_IYUV(__global const uchar* srcptr, int src_step, int src_offset,
 
377
                                __global uchar* dstptr, int dst_step, int dt_offset,
 
378
                                int rows, int cols)
 
379
{
 
380
    int x = get_global_id(0);
 
381
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
382
 
 
383
    if (x < cols / 2)
 
384
    {
 
385
        #pragma unroll
 
386
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
387
        {
 
388
            if (y < rows / 2 )
 
389
            {
 
390
                __global const uchar* ysrc = srcptr + mad24(y << 1, src_step, (x << 1) + src_offset);
 
391
                __global uchar*       dst1 = dstptr + mad24(y << 1, dst_step, x * (dcn<<1) + dt_offset);
 
392
                __global uchar*       dst2 = dst1 + dst_step;
 
393
 
 
394
                float Y1 = ysrc[0];
 
395
                float Y2 = ysrc[1];
 
396
                float Y3 = ysrc[src_step];
 
397
                float Y4 = ysrc[src_step + 1];
 
398
 
 
399
#ifdef SRC_CONT
 
400
                __global const uchar* uvsrc = srcptr + mad24(rows, src_step, src_offset);
 
401
                int u_ind = mad24(y, cols >> 1, x);
 
402
                float uv[2] = { ((float)uvsrc[u_ind]) - HALF_MAX, ((float)uvsrc[u_ind + ((rows * cols) >> 2)]) - HALF_MAX };
 
403
#else
 
404
                int vsteps[2] = { cols >> 1, src_step - (cols >> 1)};
 
405
                __global const uchar* usrc = srcptr + mad24(rows + (y>>1), src_step, src_offset + (y%2)*(cols >> 1) + x);
 
406
                __global const uchar* vsrc = usrc + mad24(rows >> 2, src_step, rows % 4 ? vsteps[y%2] : 0);
 
407
                float uv[2] = { ((float)usrc[0]) - HALF_MAX, ((float)vsrc[0]) - HALF_MAX };
 
408
#endif
 
409
                float U = uv[uidx];
 
410
                float V = uv[1-uidx];
 
411
 
 
412
                __constant float* coeffs = c_YUV2RGBCoeffs_420;
 
413
                float ruv = fma(coeffs[4], V, 0.5f);
 
414
                float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f));
 
415
                float buv = fma(coeffs[1], U, 0.5f);
 
416
 
 
417
                Y1 = max(0.f, Y1 - 16.f) * coeffs[0];
 
418
                dst1[2 - bidx] = convert_uchar_sat(Y1 + ruv);
 
419
                dst1[1]        = convert_uchar_sat(Y1 + guv);
 
420
                dst1[bidx]     = convert_uchar_sat(Y1 + buv);
 
421
#if dcn == 4
 
422
                dst1[3]        = 255;
 
423
#endif
 
424
 
 
425
                Y2 = max(0.f, Y2 - 16.f) * coeffs[0];
 
426
                dst1[dcn + 2 - bidx] = convert_uchar_sat(Y2 + ruv);
 
427
                dst1[dcn + 1]        = convert_uchar_sat(Y2 + guv);
 
428
                dst1[dcn + bidx]     = convert_uchar_sat(Y2 + buv);
 
429
#if dcn == 4
 
430
                dst1[7]        = 255;
 
431
#endif
 
432
 
 
433
                Y3 = max(0.f, Y3 - 16.f) * coeffs[0];
 
434
                dst2[2 - bidx] = convert_uchar_sat(Y3 + ruv);
 
435
                dst2[1]        = convert_uchar_sat(Y3 + guv);
 
436
                dst2[bidx]     = convert_uchar_sat(Y3 + buv);
 
437
#if dcn == 4
 
438
                dst2[3]        = 255;
 
439
#endif
 
440
 
 
441
                Y4 = max(0.f, Y4 - 16.f) * coeffs[0];
 
442
                dst2[dcn + 2 - bidx] = convert_uchar_sat(Y4 + ruv);
 
443
                dst2[dcn + 1]        = convert_uchar_sat(Y4 + guv);
 
444
                dst2[dcn + bidx]     = convert_uchar_sat(Y4 + buv);
 
445
#if dcn == 4
 
446
                dst2[7]        = 255;
 
447
#endif
 
448
            }
 
449
            ++y;
 
450
        }
 
451
    }
 
452
}
 
453
 
 
454
__constant float c_RGB2YUVCoeffs_420[8] = { 0.256999969f, 0.50399971f, 0.09799957f, -0.1479988098f, -0.2909994125f,
 
455
                                            0.438999176f, -0.3679990768f, -0.0709991455f };
 
456
 
 
457
__kernel void RGB2YUV_YV12_IYUV(__global const uchar* srcptr, int src_step, int src_offset,
 
458
                                __global uchar* dstptr, int dst_step, int dst_offset,
 
459
                                int rows, int cols)
 
460
{
 
461
    int x = get_global_id(0) * PIX_PER_WI_X;
 
462
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
463
 
 
464
    if (x < cols/2)
 
465
    {
 
466
        int src_index  = mad24(y << 1, src_step, mad24(x << 1, scn, src_offset));
 
467
        int ydst_index = mad24(y << 1, dst_step, (x << 1) + dst_offset);
 
468
        int y_rows = rows / 3 * 2;
 
469
        int vsteps[2] = { cols >> 1, dst_step - (cols >> 1)};
 
470
        __constant float* coeffs = c_RGB2YUVCoeffs_420;
 
471
 
 
472
        #pragma unroll
 
473
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
474
        {
 
475
            if (y < rows / 3)
 
476
            {
 
477
                __global const uchar* src1 = srcptr + src_index;
 
478
                __global const uchar* src2 = src1 + src_step;
 
479
                __global uchar* ydst1 = dstptr + ydst_index;
 
480
                __global uchar* ydst2 = ydst1 + dst_step;
 
481
 
 
482
                __global uchar* udst = dstptr + mad24(y_rows + (y>>1), dst_step, dst_offset + (y%2)*(cols >> 1) + x);
 
483
                __global uchar* vdst = udst + mad24(y_rows >> 2, dst_step, y_rows % 4 ? vsteps[y%2] : 0);
 
484
 
 
485
#if PIX_PER_WI_X == 2
 
486
                int s11 = *((__global const int*) src1);
 
487
                int s12 = *((__global const int*) src1 + 1);
 
488
                int s13 = *((__global const int*) src1 + 2);
 
489
#if scn == 4
 
490
                int s14 = *((__global const int*) src1 + 3);
 
491
#endif
 
492
                int s21 = *((__global const int*) src2);
 
493
                int s22 = *((__global const int*) src2 + 1);
 
494
                int s23 = *((__global const int*) src2 + 2);
 
495
#if scn == 4
 
496
                int s24 = *((__global const int*) src2 + 3);
 
497
#endif
 
498
                float src_pix1[scn * 4], src_pix2[scn * 4];
 
499
 
 
500
                *((float4*) src_pix1)     = convert_float4(as_uchar4(s11));
 
501
                *((float4*) src_pix1 + 1) = convert_float4(as_uchar4(s12));
 
502
                *((float4*) src_pix1 + 2) = convert_float4(as_uchar4(s13));
 
503
#if scn == 4
 
504
                *((float4*) src_pix1 + 3) = convert_float4(as_uchar4(s14));
 
505
#endif
 
506
                *((float4*) src_pix2)     = convert_float4(as_uchar4(s21));
 
507
                *((float4*) src_pix2 + 1) = convert_float4(as_uchar4(s22));
 
508
                *((float4*) src_pix2 + 2) = convert_float4(as_uchar4(s23));
 
509
#if scn == 4
 
510
                *((float4*) src_pix2 + 3) = convert_float4(as_uchar4(s24));
 
511
#endif
 
512
                uchar4 y1, y2;
 
513
                y1.x = convert_uchar_sat(fma(coeffs[0], src_pix1[      2-bidx], fma(coeffs[1], src_pix1[      1], fma(coeffs[2], src_pix1[      bidx], 16.5f))));
 
514
                y1.y = convert_uchar_sat(fma(coeffs[0], src_pix1[  scn+2-bidx], fma(coeffs[1], src_pix1[  scn+1], fma(coeffs[2], src_pix1[  scn+bidx], 16.5f))));
 
515
                y1.z = convert_uchar_sat(fma(coeffs[0], src_pix1[2*scn+2-bidx], fma(coeffs[1], src_pix1[2*scn+1], fma(coeffs[2], src_pix1[2*scn+bidx], 16.5f))));
 
516
                y1.w = convert_uchar_sat(fma(coeffs[0], src_pix1[3*scn+2-bidx], fma(coeffs[1], src_pix1[3*scn+1], fma(coeffs[2], src_pix1[3*scn+bidx], 16.5f))));
 
517
                y2.x = convert_uchar_sat(fma(coeffs[0], src_pix2[      2-bidx], fma(coeffs[1], src_pix2[      1], fma(coeffs[2], src_pix2[      bidx], 16.5f))));
 
518
                y2.y = convert_uchar_sat(fma(coeffs[0], src_pix2[  scn+2-bidx], fma(coeffs[1], src_pix2[  scn+1], fma(coeffs[2], src_pix2[  scn+bidx], 16.5f))));
 
519
                y2.z = convert_uchar_sat(fma(coeffs[0], src_pix2[2*scn+2-bidx], fma(coeffs[1], src_pix2[2*scn+1], fma(coeffs[2], src_pix2[2*scn+bidx], 16.5f))));
 
520
                y2.w = convert_uchar_sat(fma(coeffs[0], src_pix2[3*scn+2-bidx], fma(coeffs[1], src_pix2[3*scn+1], fma(coeffs[2], src_pix2[3*scn+bidx], 16.5f))));
 
521
 
 
522
                *((__global int*) ydst1) = as_int(y1);
 
523
                *((__global int*) ydst2) = as_int(y2);
 
524
 
 
525
                float uv[4] = { fma(coeffs[3], src_pix1[      2-bidx], fma(coeffs[4], src_pix1[      1], fma(coeffs[5], src_pix1[      bidx], 128.5f))),
 
526
                                fma(coeffs[5], src_pix1[      2-bidx], fma(coeffs[6], src_pix1[      1], fma(coeffs[7], src_pix1[      bidx], 128.5f))),
 
527
                                fma(coeffs[3], src_pix1[2*scn+2-bidx], fma(coeffs[4], src_pix1[2*scn+1], fma(coeffs[5], src_pix1[2*scn+bidx], 128.5f))),
 
528
                                fma(coeffs[5], src_pix1[2*scn+2-bidx], fma(coeffs[6], src_pix1[2*scn+1], fma(coeffs[7], src_pix1[2*scn+bidx], 128.5f))) };
 
529
 
 
530
                udst[0] = convert_uchar_sat(uv[uidx]    );
 
531
                vdst[0] = convert_uchar_sat(uv[1 - uidx]);
 
532
                udst[1] = convert_uchar_sat(uv[2 + uidx]);
 
533
                vdst[1] = convert_uchar_sat(uv[3 - uidx]);
 
534
#else
 
535
                float4 src_pix1 = convert_float4(vload4(0, src1));
 
536
                float4 src_pix2 = convert_float4(vload4(0, src1+scn));
 
537
                float4 src_pix3 = convert_float4(vload4(0, src2));
 
538
                float4 src_pix4 = convert_float4(vload4(0, src2+scn));
 
539
 
 
540
                ydst1[0] = convert_uchar_sat(fma(coeffs[0], src_pix1.R_COMP, fma(coeffs[1], src_pix1.G_COMP, fma(coeffs[2], src_pix1.B_COMP, 16.5f))));
 
541
                ydst1[1] = convert_uchar_sat(fma(coeffs[0], src_pix2.R_COMP, fma(coeffs[1], src_pix2.G_COMP, fma(coeffs[2], src_pix2.B_COMP, 16.5f))));
 
542
                ydst2[0] = convert_uchar_sat(fma(coeffs[0], src_pix3.R_COMP, fma(coeffs[1], src_pix3.G_COMP, fma(coeffs[2], src_pix3.B_COMP, 16.5f))));
 
543
                ydst2[1] = convert_uchar_sat(fma(coeffs[0], src_pix4.R_COMP, fma(coeffs[1], src_pix4.G_COMP, fma(coeffs[2], src_pix4.B_COMP, 16.5f))));
 
544
 
 
545
                float uv[2] = { fma(coeffs[3], src_pix1.R_COMP, fma(coeffs[4], src_pix1.G_COMP, fma(coeffs[5], src_pix1.B_COMP, 128.5f))),
 
546
                                fma(coeffs[5], src_pix1.R_COMP, fma(coeffs[6], src_pix1.G_COMP, fma(coeffs[7], src_pix1.B_COMP, 128.5f))) };
 
547
 
 
548
                udst[0] = convert_uchar_sat(uv[uidx]  );
 
549
                vdst[0] = convert_uchar_sat(uv[1-uidx]);
 
550
#endif
 
551
                ++y;
 
552
                src_index += 2*src_step;
 
553
                ydst_index += 2*dst_step;
 
554
            }
 
555
        }
 
556
    }
 
557
}
 
558
 
 
559
__kernel void YUV2RGB_422(__global const uchar* srcptr, int src_step, int src_offset,
 
560
                          __global uchar* dstptr, int dst_step, int dst_offset,
 
561
                          int rows, int cols)
 
562
{
 
563
    int x = get_global_id(0);
 
564
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
565
 
 
566
    if (x < cols / 2)
 
567
    {
 
568
        __global const uchar* src = srcptr + mad24(y, src_step, (x << 2) + src_offset);
 
569
        __global uchar*       dst = dstptr + mad24(y, dst_step, mad24(x << 1, dcn, dst_offset));
 
570
 
 
571
        #pragma unroll
 
572
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
573
        {
 
574
            if (y < rows )
 
575
            {
 
576
                __constant float* coeffs = c_YUV2RGBCoeffs_420;
 
577
 
 
578
#ifndef USE_OPTIMIZED_LOAD
 
579
                float U = ((float) src[uidx]) - HALF_MAX;
 
580
                float V = ((float) src[(2 + uidx) % 4]) - HALF_MAX;
 
581
                float y00 = max(0.f, ((float) src[yidx]) - 16.f) * coeffs[0];
 
582
                float y01 = max(0.f, ((float) src[yidx + 2]) - 16.f) * coeffs[0];
 
583
#else
 
584
                int load_src = *((__global int*) src);
 
585
                float vec_src[4] = { load_src & 0xff, (load_src >> 8) & 0xff, (load_src >> 16) & 0xff, (load_src >> 24) & 0xff};
 
586
                float U = vec_src[uidx] - HALF_MAX;
 
587
                float V = vec_src[(2 + uidx) % 4] - HALF_MAX;
 
588
                float y00 = max(0.f, vec_src[yidx] - 16.f) * coeffs[0];
 
589
                float y01 = max(0.f, vec_src[yidx + 2] - 16.f) * coeffs[0];
 
590
#endif
 
591
 
 
592
                float ruv = fma(coeffs[4], V, 0.5f);
 
593
                float guv = fma(coeffs[3], V, fma(coeffs[2], U, 0.5f));
 
594
                float buv = fma(coeffs[1], U, 0.5f);
 
595
 
 
596
                dst[2 - bidx] = convert_uchar_sat(y00 + ruv);
 
597
                dst[1]        = convert_uchar_sat(y00 + guv);
 
598
                dst[bidx]     = convert_uchar_sat(y00 + buv);
 
599
#if dcn == 4
 
600
                dst[3]        = 255;
 
601
#endif
 
602
 
 
603
                dst[dcn + 2 - bidx] = convert_uchar_sat(y01 + ruv);
 
604
                dst[dcn + 1]        = convert_uchar_sat(y01 + guv);
 
605
                dst[dcn + bidx]     = convert_uchar_sat(y01 + buv);
 
606
#if dcn == 4
 
607
                dst[7]        = 255;
 
608
#endif
 
609
            }
 
610
            ++y;
 
611
            src += src_step;
 
612
            dst += dst_step;
 
613
        }
 
614
    }
 
615
}
 
616
 
 
617
///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
 
618
 
 
619
__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
 
620
__constant int   c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
 
621
 
 
622
__kernel void RGB2YCrCb(__global const uchar* srcptr, int src_step, int src_offset,
 
623
                        __global uchar* dstptr, int dst_step, int dt_offset,
 
624
                        int rows, int cols)
 
625
{
 
626
    int x = get_global_id(0);
 
627
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
628
 
 
629
    if (x < cols)
 
630
    {
 
631
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
632
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dt_offset));
 
633
 
 
634
        #pragma unroll
 
635
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
636
        {
 
637
            if (y < rows)
 
638
            {
 
639
                __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
 
640
                __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
 
641
                DATA_TYPE_4 src_pix = vload4(0, src);
 
642
                DATA_TYPE b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
643
 
 
644
#ifdef DEPTH_5
 
645
                __constant float * coeffs = c_RGB2YCrCbCoeffs_f;
 
646
                DATA_TYPE Y = fma(b, coeffs[2], fma(g, coeffs[1], r * coeffs[0]));
 
647
                DATA_TYPE Cr = fma(r - Y, coeffs[3], HALF_MAX);
 
648
                DATA_TYPE Cb = fma(b - Y, coeffs[4], HALF_MAX);
 
649
#else
 
650
                __constant int * coeffs = c_RGB2YCrCbCoeffs_i;
 
651
                int delta = HALF_MAX * (1 << yuv_shift);
 
652
                int Y =  CV_DESCALE(mad24(b, coeffs[2], mad24(g, coeffs[1], mul24(r, coeffs[0]))), yuv_shift);
 
653
                int Cr = CV_DESCALE(mad24(r - Y, coeffs[3], delta), yuv_shift);
 
654
                int Cb = CV_DESCALE(mad24(b - Y, coeffs[4], delta), yuv_shift);
 
655
#endif
 
656
 
 
657
                dst[0] = SAT_CAST( Y );
 
658
                dst[1] = SAT_CAST( Cr );
 
659
                dst[2] = SAT_CAST( Cb );
 
660
 
 
661
                ++y;
 
662
                dst_index += dst_step;
 
663
                src_index += src_step;
 
664
            }
 
665
        }
 
666
    }
 
667
}
 
668
 
 
669
__constant float c_YCrCb2RGBCoeffs_f[4] = { 1.403f, -0.714f, -0.344f, 1.773f };
 
670
__constant int   c_YCrCb2RGBCoeffs_i[4] = { 22987, -11698, -5636, 29049 };
 
671
 
 
672
__kernel void YCrCb2RGB(__global const uchar* src, int src_step, int src_offset,
 
673
                        __global uchar* dst, int dst_step, int dst_offset,
 
674
                        int rows, int cols)
 
675
{
 
676
    int x = get_global_id(0);
 
677
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
678
 
 
679
    if (x < cols)
 
680
    {
 
681
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
682
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
683
 
 
684
        #pragma unroll
 
685
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
686
        {
 
687
            if (y < rows)
 
688
            {
 
689
                __global const DATA_TYPE * srcptr = (__global const DATA_TYPE*)(src + src_index);
 
690
                __global DATA_TYPE * dstptr = (__global DATA_TYPE*)(dst + dst_index);
 
691
 
 
692
                DATA_TYPE_4 src_pix = vload4(0, srcptr);
 
693
                DATA_TYPE yp = src_pix.x, cr = src_pix.y, cb = src_pix.z;
 
694
 
 
695
#ifdef DEPTH_5
 
696
                __constant float * coeff = c_YCrCb2RGBCoeffs_f;
 
697
                float r = fma(coeff[0], cr - HALF_MAX, yp);
 
698
                float g = fma(coeff[1], cr - HALF_MAX, fma(coeff[2], cb - HALF_MAX, yp));
 
699
                float b = fma(coeff[3], cb - HALF_MAX, yp);
 
700
#else
 
701
                __constant int * coeff = c_YCrCb2RGBCoeffs_i;
 
702
                int r = yp + CV_DESCALE(coeff[0] * (cr - HALF_MAX), yuv_shift);
 
703
                int g = yp + CV_DESCALE(mad24(coeff[1], cr - HALF_MAX, coeff[2] * (cb - HALF_MAX)), yuv_shift);
 
704
                int b = yp + CV_DESCALE(coeff[3] * (cb - HALF_MAX), yuv_shift);
 
705
#endif
 
706
 
 
707
                dstptr[(bidx^2)] = SAT_CAST(r);
 
708
                dstptr[1] = SAT_CAST(g);
 
709
                dstptr[bidx] = SAT_CAST(b);
 
710
#if dcn == 4
 
711
                dstptr[3] = MAX_NUM;
 
712
#endif
 
713
 
 
714
                ++y;
 
715
                dst_index += dst_step;
 
716
                src_index += src_step;
 
717
            }
 
718
        }
 
719
    }
 
720
}
 
721
 
 
722
///////////////////////////////////// RGB <-> XYZ //////////////////////////////////////
 
723
 
 
724
__kernel void RGB2XYZ(__global const uchar * srcptr, int src_step, int src_offset,
 
725
                      __global uchar * dstptr, int dst_step, int dst_offset,
 
726
                      int rows, int cols, __constant COEFF_TYPE * coeffs)
 
727
{
 
728
    int dx = get_global_id(0);
 
729
    int dy = get_global_id(1) * PIX_PER_WI_Y;
 
730
 
 
731
    if (dx < cols)
 
732
    {
 
733
        int src_index = mad24(dy, src_step, mad24(dx, scnbytes, src_offset));
 
734
        int dst_index = mad24(dy, dst_step, mad24(dx, dcnbytes, dst_offset));
 
735
 
 
736
        #pragma unroll
 
737
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
738
        {
 
739
            if (dy < rows)
 
740
            {
 
741
                __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index);
 
742
                __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index);
 
743
 
 
744
                DATA_TYPE_4 src_pix = vload4(0, src);
 
745
                DATA_TYPE r = src_pix.x, g = src_pix.y, b = src_pix.z;
 
746
 
 
747
#ifdef DEPTH_5
 
748
                float x = fma(r, coeffs[0], fma(g, coeffs[1], b * coeffs[2]));
 
749
                float y = fma(r, coeffs[3], fma(g, coeffs[4], b * coeffs[5]));
 
750
                float z = fma(r, coeffs[6], fma(g, coeffs[7], b * coeffs[8]));
 
751
#else
 
752
                int x = CV_DESCALE(mad24(r, coeffs[0], mad24(g, coeffs[1], b * coeffs[2])), xyz_shift);
 
753
                int y = CV_DESCALE(mad24(r, coeffs[3], mad24(g, coeffs[4], b * coeffs[5])), xyz_shift);
 
754
                int z = CV_DESCALE(mad24(r, coeffs[6], mad24(g, coeffs[7], b * coeffs[8])), xyz_shift);
 
755
#endif
 
756
                dst[0] = SAT_CAST(x);
 
757
                dst[1] = SAT_CAST(y);
 
758
                dst[2] = SAT_CAST(z);
 
759
 
 
760
                ++dy;
 
761
                dst_index += dst_step;
 
762
                src_index += src_step;
 
763
            }
 
764
        }
 
765
    }
 
766
}
 
767
 
 
768
__kernel void XYZ2RGB(__global const uchar * srcptr, int src_step, int src_offset,
 
769
                      __global uchar * dstptr, int dst_step, int dst_offset,
 
770
                      int rows, int cols, __constant COEFF_TYPE * coeffs)
 
771
{
 
772
    int dx = get_global_id(0);
 
773
    int dy = get_global_id(1) * PIX_PER_WI_Y;
 
774
 
 
775
    if (dx < cols)
 
776
    {
 
777
        int src_index = mad24(dy, src_step, mad24(dx, scnbytes, src_offset));
 
778
        int dst_index = mad24(dy, dst_step, mad24(dx, dcnbytes, dst_offset));
 
779
 
 
780
        #pragma unroll
 
781
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
782
        {
 
783
            if (dy < rows)
 
784
            {
 
785
                __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index);
 
786
                __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index);
 
787
 
 
788
                DATA_TYPE_4 src_pix = vload4(0, src);
 
789
                DATA_TYPE x = src_pix.x, y = src_pix.y, z = src_pix.z;
 
790
 
 
791
#ifdef DEPTH_5
 
792
                float b = fma(x, coeffs[0], fma(y, coeffs[1], z * coeffs[2]));
 
793
                float g = fma(x, coeffs[3], fma(y, coeffs[4], z * coeffs[5]));
 
794
                float r = fma(x, coeffs[6], fma(y, coeffs[7], z * coeffs[8]));
 
795
#else
 
796
                int b = CV_DESCALE(mad24(x, coeffs[0], mad24(y, coeffs[1], z * coeffs[2])), xyz_shift);
 
797
                int g = CV_DESCALE(mad24(x, coeffs[3], mad24(y, coeffs[4], z * coeffs[5])), xyz_shift);
 
798
                int r = CV_DESCALE(mad24(x, coeffs[6], mad24(y, coeffs[7], z * coeffs[8])), xyz_shift);
 
799
#endif
 
800
 
 
801
                DATA_TYPE dst0 = SAT_CAST(b);
 
802
                DATA_TYPE dst1 = SAT_CAST(g);
 
803
                DATA_TYPE dst2 = SAT_CAST(r);
 
804
#if dcn == 3 || defined DEPTH_5
 
805
                dst[0] = dst0;
 
806
                dst[1] = dst1;
 
807
                dst[2] = dst2;
 
808
#if dcn == 4
 
809
                dst[3] = MAX_NUM;
 
810
#endif
 
811
#else
 
812
                *(__global DATA_TYPE_4 *)dst = (DATA_TYPE_4)(dst0, dst1, dst2, MAX_NUM);
 
813
#endif
 
814
 
 
815
                ++dy;
 
816
                dst_index += dst_step;
 
817
                src_index += src_step;
 
818
            }
 
819
        }
 
820
    }
 
821
}
 
822
 
 
823
///////////////////////////////////// RGB[A] <-> BGR[A] //////////////////////////////////////
 
824
 
 
825
__kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset,
 
826
                  __global uchar* dstptr, int dst_step, int dst_offset,
 
827
                  int rows, int cols)
 
828
{
 
829
    int x = get_global_id(0);
 
830
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
831
 
 
832
    if (x < cols)
 
833
    {
 
834
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
835
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
836
 
 
837
        #pragma unroll
 
838
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
839
        {
 
840
            if (y < rows)
 
841
            {
 
842
                __global const DATA_TYPE * src = (__global const DATA_TYPE *)(srcptr + src_index);
 
843
                __global DATA_TYPE * dst = (__global DATA_TYPE *)(dstptr + dst_index);
 
844
                DATA_TYPE_4 src_pix = vload4(0, src);
 
845
 
 
846
#ifdef REVERSE
 
847
                dst[0] = src_pix.z;
 
848
                dst[1] = src_pix.y;
 
849
                dst[2] = src_pix.x;
 
850
#else
 
851
                dst[0] = src_pix.x;
 
852
                dst[1] = src_pix.y;
 
853
                dst[2] = src_pix.z;
 
854
#endif
 
855
 
 
856
#if dcn == 4
 
857
#if scn == 3
 
858
                dst[3] = MAX_NUM;
 
859
#else
 
860
                dst[3] = src[3];
 
861
#endif
 
862
#endif
 
863
 
 
864
                ++y;
 
865
                dst_index += dst_step;
 
866
                src_index += src_step;
 
867
            }
 
868
        }
 
869
    }
 
870
}
 
871
 
 
872
///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
 
873
 
 
874
__kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset,
 
875
                         __global uchar* dst, int dst_step, int dst_offset,
 
876
                         int rows, int cols)
 
877
{
 
878
    int x = get_global_id(0);
 
879
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
880
 
 
881
    if (x < cols)
 
882
    {
 
883
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
884
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
885
 
 
886
        #pragma unroll
 
887
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
888
        {
 
889
            if (y < rows)
 
890
            {
 
891
                ushort t = *((__global const ushort*)(src + src_index));
 
892
 
 
893
#if greenbits == 6
 
894
                dst[dst_index + bidx] = (uchar)(t << 3);
 
895
                dst[dst_index + 1] = (uchar)((t >> 3) & ~3);
 
896
                dst[dst_index + (bidx^2)] = (uchar)((t >> 8) & ~7);
 
897
#else
 
898
                dst[dst_index + bidx] = (uchar)(t << 3);
 
899
                dst[dst_index + 1] = (uchar)((t >> 2) & ~7);
 
900
                dst[dst_index + (bidx^2)] = (uchar)((t >> 7) & ~7);
 
901
#endif
 
902
 
 
903
#if dcn == 4
 
904
#if greenbits == 6
 
905
                dst[dst_index + 3] = 255;
 
906
#else
 
907
                dst[dst_index + 3] = t & 0x8000 ? 255 : 0;
 
908
#endif
 
909
#endif
 
910
 
 
911
                ++y;
 
912
                dst_index += dst_step;
 
913
                src_index += src_step;
 
914
            }
 
915
        }
 
916
    }
 
917
}
 
918
 
 
919
__kernel void RGB2RGB5x5(__global const uchar* src, int src_step, int src_offset,
 
920
                         __global uchar* dst, int dst_step, int dst_offset,
 
921
                         int rows, int cols)
 
922
{
 
923
    int x = get_global_id(0);
 
924
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
925
 
 
926
    if (x < cols)
 
927
    {
 
928
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
929
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
930
 
 
931
        #pragma unroll
 
932
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
933
        {
 
934
            if (y < rows)
 
935
            {
 
936
                uchar4 src_pix = vload4(0, src + src_index);
 
937
 
 
938
#if greenbits == 6
 
939
                    *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~3) << 3)|((src_pix.R_COMP&~7) << 8));
 
940
#elif scn == 3
 
941
                    *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)|((src_pix.R_COMP&~7) << 7));
 
942
#else
 
943
                    *((__global ushort*)(dst + dst_index)) = (ushort)((src_pix.B_COMP >> 3)|((src_pix.G_COMP&~7) << 2)|
 
944
                        ((src_pix.R_COMP&~7) << 7)|(src_pix.w ? 0x8000 : 0));
 
945
#endif
 
946
 
 
947
                ++y;
 
948
                dst_index += dst_step;
 
949
                src_index += src_step;
 
950
            }
 
951
        }
 
952
    }
 
953
}
 
954
 
 
955
///////////////////////////////////// RGB5x5 <-> Gray //////////////////////////////////////
 
956
 
 
957
__kernel void BGR5x52Gray(__global const uchar* src, int src_step, int src_offset,
 
958
                          __global uchar* dst, int dst_step, int dst_offset,
 
959
                          int rows, int cols)
 
960
{
 
961
    int x = get_global_id(0);
 
962
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
963
 
 
964
    if (x < cols)
 
965
    {
 
966
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
967
        int dst_index = mad24(y, dst_step, dst_offset + x);
 
968
 
 
969
        #pragma unroll
 
970
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
971
        {
 
972
            if (y < rows)
 
973
            {
 
974
                int t = *((__global const ushort*)(src + src_index));
 
975
 
 
976
#if greenbits == 6
 
977
                dst[dst_index] = (uchar)CV_DESCALE(mad24((t << 3) & 0xf8, B2Y, mad24((t >> 3) & 0xfc, G2Y, ((t >> 8) & 0xf8) * R2Y)), yuv_shift);
 
978
#else
 
979
                dst[dst_index] = (uchar)CV_DESCALE(mad24((t << 3) & 0xf8, B2Y, mad24((t >> 2) & 0xf8, G2Y, ((t >> 7) & 0xf8) * R2Y)), yuv_shift);
 
980
#endif
 
981
                ++y;
 
982
                dst_index += dst_step;
 
983
                src_index += src_step;
 
984
            }
 
985
        }
 
986
    }
 
987
}
 
988
 
 
989
__kernel void Gray2BGR5x5(__global const uchar* src, int src_step, int src_offset,
 
990
                          __global uchar* dst, int dst_step, int dst_offset,
 
991
                          int rows, int cols)
 
992
{
 
993
    int x = get_global_id(0);
 
994
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
995
 
 
996
    if (x < cols)
 
997
    {
 
998
        int src_index = mad24(y, src_step, src_offset + x);
 
999
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1000
 
 
1001
        #pragma unroll
 
1002
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1003
        {
 
1004
            if (y < rows)
 
1005
            {
 
1006
                int t = src[src_index];
 
1007
 
 
1008
#if greenbits == 6
 
1009
                *((__global ushort*)(dst + dst_index)) = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8));
 
1010
#else
 
1011
                t >>= 3;
 
1012
                *((__global ushort*)(dst + dst_index)) = (ushort)(t|(t << 5)|(t << 10));
 
1013
#endif
 
1014
                ++y;
 
1015
                dst_index += dst_step;
 
1016
                src_index += src_step;
 
1017
            }
 
1018
        }
 
1019
    }
 
1020
}
 
1021
 
 
1022
//////////////////////////////////// RGB <-> HSV //////////////////////////////////////
 
1023
 
 
1024
__constant int sector_data[][3] = { { 1, 3, 0 },
 
1025
                                    { 1, 0, 2 },
 
1026
                                    { 3, 0, 1 },
 
1027
                                    { 0, 2, 1 },
 
1028
                                    { 0, 1, 3 },
 
1029
                                    { 2, 1, 0 } };
 
1030
 
 
1031
#ifdef DEPTH_0
 
1032
 
 
1033
__kernel void RGB2HSV(__global const uchar* src, int src_step, int src_offset,
 
1034
                      __global uchar* dst, int dst_step, int dst_offset,
 
1035
                      int rows, int cols,
 
1036
                      __constant int * sdiv_table, __constant int * hdiv_table)
 
1037
{
 
1038
    int x = get_global_id(0);
 
1039
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1040
 
 
1041
    if (x < cols)
 
1042
    {
 
1043
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1044
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1045
 
 
1046
        #pragma unroll
 
1047
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1048
        {
 
1049
            if (y < rows)
 
1050
            {
 
1051
                uchar4 src_pix = vload4(0, src + src_index);
 
1052
 
 
1053
                int b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
1054
                int h, s, v = b;
 
1055
                int vmin = b, diff;
 
1056
                int vr, vg;
 
1057
 
 
1058
                v = max(v, g);
 
1059
                v = max(v, r);
 
1060
                vmin = min(vmin, g);
 
1061
                vmin = min(vmin, r);
 
1062
 
 
1063
                diff = v - vmin;
 
1064
                vr = v == r ? -1 : 0;
 
1065
                vg = v == g ? -1 : 0;
 
1066
 
 
1067
                s = mad24(diff, sdiv_table[v], (1 << (hsv_shift-1))) >> hsv_shift;
 
1068
                h = (vr & (g - b)) +
 
1069
                    (~vr & ((vg & mad24(diff, 2, b - r)) + ((~vg) & mad24(4, diff, r - g))));
 
1070
                h = mad24(h, hdiv_table[diff], (1 << (hsv_shift-1))) >> hsv_shift;
 
1071
                h += h < 0 ? hrange : 0;
 
1072
 
 
1073
                dst[dst_index] = convert_uchar_sat_rte(h);
 
1074
                dst[dst_index + 1] = (uchar)s;
 
1075
                dst[dst_index + 2] = (uchar)v;
 
1076
 
 
1077
                ++y;
 
1078
                dst_index += dst_step;
 
1079
                src_index += src_step;
 
1080
            }
 
1081
        }
 
1082
    }
 
1083
}
 
1084
 
 
1085
__kernel void HSV2RGB(__global const uchar* src, int src_step, int src_offset,
 
1086
                      __global uchar* dst, int dst_step, int dst_offset,
 
1087
                      int rows, int cols)
 
1088
{
 
1089
    int x = get_global_id(0);
 
1090
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1091
 
 
1092
    if (x < cols)
 
1093
    {
 
1094
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1095
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1096
 
 
1097
        #pragma unroll
 
1098
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1099
        {
 
1100
            if (y < rows)
 
1101
            {
 
1102
                uchar4 src_pix = vload4(0, src + src_index);
 
1103
 
 
1104
                float h = src_pix.x, s = src_pix.y*(1/255.f), v = src_pix.z*(1/255.f);
 
1105
                float b, g, r;
 
1106
 
 
1107
                if (s != 0)
 
1108
                {
 
1109
                    float tab[4];
 
1110
                    int sector;
 
1111
                    h *= hscale;
 
1112
                    if( h < 0 )
 
1113
                        do h += 6; while( h < 0 );
 
1114
                    else if( h >= 6 )
 
1115
                        do h -= 6; while( h >= 6 );
 
1116
                    sector = convert_int_sat_rtn(h);
 
1117
                    h -= sector;
 
1118
                    if( (unsigned)sector >= 6u )
 
1119
                    {
 
1120
                        sector = 0;
 
1121
                        h = 0.f;
 
1122
                    }
 
1123
 
 
1124
                    tab[0] = v;
 
1125
                    tab[1] = v*(1.f - s);
 
1126
                    tab[2] = v*(1.f - s*h);
 
1127
                    tab[3] = v*(1.f - s*(1.f - h));
 
1128
 
 
1129
                    b = tab[sector_data[sector][0]];
 
1130
                    g = tab[sector_data[sector][1]];
 
1131
                    r = tab[sector_data[sector][2]];
 
1132
                }
 
1133
                else
 
1134
                    b = g = r = v;
 
1135
 
 
1136
                dst[dst_index + bidx] = convert_uchar_sat_rte(b*255.f);
 
1137
                dst[dst_index + 1] = convert_uchar_sat_rte(g*255.f);
 
1138
                dst[dst_index + (bidx^2)] = convert_uchar_sat_rte(r*255.f);
 
1139
#if dcn == 4
 
1140
                dst[dst_index + 3] = MAX_NUM;
 
1141
#endif
 
1142
 
 
1143
                ++y;
 
1144
                dst_index += dst_step;
 
1145
                src_index += src_step;
 
1146
            }
 
1147
        }
 
1148
    }
 
1149
}
 
1150
 
 
1151
#elif defined DEPTH_5
 
1152
 
 
1153
__kernel void RGB2HSV(__global const uchar* srcptr, int src_step, int src_offset,
 
1154
                      __global uchar* dstptr, int dst_step, int dst_offset,
 
1155
                      int rows, int cols)
 
1156
{
 
1157
    int x = get_global_id(0);
 
1158
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1159
 
 
1160
    if (x < cols)
 
1161
    {
 
1162
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1163
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1164
 
 
1165
        #pragma unroll
 
1166
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1167
        {
 
1168
            if (y < rows)
 
1169
            {
 
1170
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1171
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1172
                float4 src_pix = vload4(0, src);
 
1173
 
 
1174
                float b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
1175
                float h, s, v;
 
1176
 
 
1177
                float vmin, diff;
 
1178
 
 
1179
                v = vmin = r;
 
1180
                if( v < g ) v = g;
 
1181
                if( v < b ) v = b;
 
1182
                if( vmin > g ) vmin = g;
 
1183
                if( vmin > b ) vmin = b;
 
1184
 
 
1185
                diff = v - vmin;
 
1186
                s = diff/(float)(fabs(v) + FLT_EPSILON);
 
1187
                diff = (float)(60.f/(diff + FLT_EPSILON));
 
1188
                if( v == r )
 
1189
                    h = (g - b)*diff;
 
1190
                else if( v == g )
 
1191
                    h = fma(b - r, diff, 120.f);
 
1192
                else
 
1193
                    h = fma(r - g, diff, 240.f);
 
1194
 
 
1195
                if( h < 0 )
 
1196
                    h += 360.f;
 
1197
 
 
1198
                dst[0] = h*hscale;
 
1199
                dst[1] = s;
 
1200
                dst[2] = v;
 
1201
 
 
1202
                ++y;
 
1203
                dst_index += dst_step;
 
1204
                src_index += src_step;
 
1205
            }
 
1206
        }
 
1207
    }
 
1208
}
 
1209
 
 
1210
__kernel void HSV2RGB(__global const uchar* srcptr, int src_step, int src_offset,
 
1211
                      __global uchar* dstptr, int dst_step, int dst_offset,
 
1212
                      int rows, int cols)
 
1213
{
 
1214
    int x = get_global_id(0);
 
1215
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1216
 
 
1217
    if (x < cols)
 
1218
    {
 
1219
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1220
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1221
 
 
1222
        #pragma unroll
 
1223
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1224
        {
 
1225
            if (y < rows)
 
1226
            {
 
1227
 
 
1228
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1229
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1230
                float4 src_pix = vload4(0, src);
 
1231
 
 
1232
                float h = src_pix.x, s = src_pix.y, v = src_pix.z;
 
1233
                float b, g, r;
 
1234
 
 
1235
                if (s != 0)
 
1236
                {
 
1237
                    float tab[4];
 
1238
                    int sector;
 
1239
                    h *= hscale;
 
1240
                    if(h < 0)
 
1241
                        do h += 6; while (h < 0);
 
1242
                    else if (h >= 6)
 
1243
                        do h -= 6; while (h >= 6);
 
1244
                    sector = convert_int_sat_rtn(h);
 
1245
                    h -= sector;
 
1246
                    if ((unsigned)sector >= 6u)
 
1247
                    {
 
1248
                        sector = 0;
 
1249
                        h = 0.f;
 
1250
                    }
 
1251
 
 
1252
                    tab[0] = v;
 
1253
                    tab[1] = v*(1.f - s);
 
1254
                    tab[2] = v*(1.f - s*h);
 
1255
                    tab[3] = v*(1.f - s*(1.f - h));
 
1256
 
 
1257
                    b = tab[sector_data[sector][0]];
 
1258
                    g = tab[sector_data[sector][1]];
 
1259
                    r = tab[sector_data[sector][2]];
 
1260
                }
 
1261
                else
 
1262
                    b = g = r = v;
 
1263
 
 
1264
                dst[bidx] = b;
 
1265
                dst[1] = g;
 
1266
                dst[bidx^2] = r;
 
1267
#if dcn == 4
 
1268
                dst[3] = MAX_NUM;
 
1269
#endif
 
1270
 
 
1271
                ++y;
 
1272
                dst_index += dst_step;
 
1273
                src_index += src_step;
 
1274
            }
 
1275
        }
 
1276
    }
 
1277
}
 
1278
 
 
1279
#endif
 
1280
 
 
1281
///////////////////////////////////// RGB <-> HLS //////////////////////////////////////
 
1282
 
 
1283
#ifdef DEPTH_0
 
1284
 
 
1285
__kernel void RGB2HLS(__global const uchar* src, int src_step, int src_offset,
 
1286
                      __global uchar* dst, int dst_step, int dst_offset,
 
1287
                      int rows, int cols)
 
1288
{
 
1289
    int x = get_global_id(0);
 
1290
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1291
 
 
1292
    if (x < cols)
 
1293
    {
 
1294
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1295
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1296
 
 
1297
        #pragma unroll
 
1298
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1299
        {
 
1300
            if (y < rows)
 
1301
            {
 
1302
                uchar4 src_pix = vload4(0, src + src_index);
 
1303
 
 
1304
                float b = src_pix.B_COMP*(1/255.f), g = src_pix.G_COMP*(1/255.f), r = src_pix.R_COMP*(1/255.f);
 
1305
                float h = 0.f, s = 0.f, l;
 
1306
                float vmin, vmax, diff;
 
1307
 
 
1308
                vmax = vmin = r;
 
1309
                if (vmax < g) vmax = g;
 
1310
                if (vmax < b) vmax = b;
 
1311
                if (vmin > g) vmin = g;
 
1312
                if (vmin > b) vmin = b;
 
1313
 
 
1314
                diff = vmax - vmin;
 
1315
                l = (vmax + vmin)*0.5f;
 
1316
 
 
1317
                if (diff > FLT_EPSILON)
 
1318
                {
 
1319
                    s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin);
 
1320
                    diff = 60.f/diff;
 
1321
 
 
1322
                    if( vmax == r )
 
1323
                        h = (g - b)*diff;
 
1324
                    else if( vmax == g )
 
1325
                        h = fma(b - r, diff, 120.f);
 
1326
                    else
 
1327
                        h = fma(r - g, diff, 240.f);
 
1328
 
 
1329
                    if( h < 0.f )
 
1330
                        h += 360.f;
 
1331
                }
 
1332
 
 
1333
                dst[dst_index] = convert_uchar_sat_rte(h*hscale);
 
1334
                dst[dst_index + 1] = convert_uchar_sat_rte(l*255.f);
 
1335
                dst[dst_index + 2] = convert_uchar_sat_rte(s*255.f);
 
1336
 
 
1337
                ++y;
 
1338
                dst_index += dst_step;
 
1339
                src_index += src_step;
 
1340
            }
 
1341
        }
 
1342
    }
 
1343
}
 
1344
 
 
1345
__kernel void HLS2RGB(__global const uchar* src, int src_step, int src_offset,
 
1346
                      __global uchar* dst, int dst_step, int dst_offset,
 
1347
                      int rows, int cols)
 
1348
{
 
1349
    int x = get_global_id(0);
 
1350
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1351
 
 
1352
    if (x < cols)
 
1353
    {
 
1354
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1355
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1356
 
 
1357
        #pragma unroll
 
1358
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1359
        {
 
1360
            if (y < rows)
 
1361
            {
 
1362
                uchar4 src_pix = vload4(0, src + src_index);
 
1363
 
 
1364
                float h = src_pix.x, l = src_pix.y*(1.f/255.f), s = src_pix.z*(1.f/255.f);
 
1365
                float b, g, r;
 
1366
 
 
1367
                if (s != 0)
 
1368
                {
 
1369
                    float tab[4];
 
1370
 
 
1371
                    float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s;
 
1372
                    float p1 = 2*l - p2;
 
1373
 
 
1374
                    h *= hscale;
 
1375
                    if( h < 0 )
 
1376
                        do h += 6; while( h < 0 );
 
1377
                    else if( h >= 6 )
 
1378
                        do h -= 6; while( h >= 6 );
 
1379
 
 
1380
                    int sector = convert_int_sat_rtn(h);
 
1381
                    h -= sector;
 
1382
 
 
1383
                    tab[0] = p2;
 
1384
                    tab[1] = p1;
 
1385
                    tab[2] = fma(p2 - p1, 1-h, p1);
 
1386
                    tab[3] = fma(p2 - p1, h, p1);
 
1387
 
 
1388
                    b = tab[sector_data[sector][0]];
 
1389
                    g = tab[sector_data[sector][1]];
 
1390
                    r = tab[sector_data[sector][2]];
 
1391
                }
 
1392
                else
 
1393
                    b = g = r = l;
 
1394
 
 
1395
                dst[dst_index + bidx] = convert_uchar_sat_rte(b*255.f);
 
1396
                dst[dst_index + 1] = convert_uchar_sat_rte(g*255.f);
 
1397
                dst[dst_index + (bidx^2)] = convert_uchar_sat_rte(r*255.f);
 
1398
#if dcn == 4
 
1399
                dst[dst_index + 3] = MAX_NUM;
 
1400
#endif
 
1401
 
 
1402
                ++y;
 
1403
                dst_index += dst_step;
 
1404
                src_index += src_step;
 
1405
            }
 
1406
        }
 
1407
    }
 
1408
}
 
1409
 
 
1410
#elif defined DEPTH_5
 
1411
 
 
1412
__kernel void RGB2HLS(__global const uchar* srcptr, int src_step, int src_offset,
 
1413
                      __global uchar* dstptr, int dst_step, int dst_offset,
 
1414
                      int rows, int cols)
 
1415
{
 
1416
    int x = get_global_id(0);
 
1417
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1418
 
 
1419
    if (x < cols)
 
1420
    {
 
1421
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1422
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1423
 
 
1424
        #pragma unroll
 
1425
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1426
        {
 
1427
            if (y < rows)
 
1428
            {
 
1429
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1430
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1431
                float4 src_pix = vload4(0, src);
 
1432
 
 
1433
                float b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
1434
                float h = 0.f, s = 0.f, l;
 
1435
                float vmin, vmax, diff;
 
1436
 
 
1437
                vmax = vmin = r;
 
1438
                if (vmax < g) vmax = g;
 
1439
                if (vmax < b) vmax = b;
 
1440
                if (vmin > g) vmin = g;
 
1441
                if (vmin > b) vmin = b;
 
1442
 
 
1443
                diff = vmax - vmin;
 
1444
                l = (vmax + vmin)*0.5f;
 
1445
 
 
1446
                if (diff > FLT_EPSILON)
 
1447
                {
 
1448
                    s = l < 0.5f ? diff/(vmax + vmin) : diff/(2 - vmax - vmin);
 
1449
                    diff = 60.f/diff;
 
1450
 
 
1451
                    if( vmax == r )
 
1452
                        h = (g - b)*diff;
 
1453
                    else if( vmax == g )
 
1454
                        h = fma(b - r, diff, 120.f);
 
1455
                    else
 
1456
                        h = fma(r - g, diff, 240.f);
 
1457
 
 
1458
                    if( h < 0.f ) h += 360.f;
 
1459
                }
 
1460
 
 
1461
                dst[0] = h*hscale;
 
1462
                dst[1] = l;
 
1463
                dst[2] = s;
 
1464
 
 
1465
                ++y;
 
1466
                dst_index += dst_step;
 
1467
                src_index += src_step;
 
1468
            }
 
1469
        }
 
1470
    }
 
1471
}
 
1472
 
 
1473
__kernel void HLS2RGB(__global const uchar* srcptr, int src_step, int src_offset,
 
1474
                      __global uchar* dstptr, int dst_step, int dst_offset,
 
1475
                      int rows, int cols)
 
1476
{
 
1477
    int x = get_global_id(0);
 
1478
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1479
 
 
1480
    if (x < cols)
 
1481
    {
 
1482
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1483
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1484
 
 
1485
        #pragma unroll
 
1486
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1487
        {
 
1488
            if (y < rows)
 
1489
            {
 
1490
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1491
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1492
                float4 src_pix = vload4(0, src);
 
1493
 
 
1494
                float h = src_pix.x, l = src_pix.y, s = src_pix.z;
 
1495
                float b, g, r;
 
1496
 
 
1497
                if (s != 0)
 
1498
                {
 
1499
                    float tab[4];
 
1500
                    int sector;
 
1501
 
 
1502
                    float p2 = l <= 0.5f ? l*(1 + s) : l + s - l*s;
 
1503
                    float p1 = 2*l - p2;
 
1504
 
 
1505
                    h *= hscale;
 
1506
                    if( h < 0 )
 
1507
                        do h += 6; while( h < 0 );
 
1508
                    else if( h >= 6 )
 
1509
                        do h -= 6; while( h >= 6 );
 
1510
 
 
1511
                    sector = convert_int_sat_rtn(h);
 
1512
                    h -= sector;
 
1513
 
 
1514
                    tab[0] = p2;
 
1515
                    tab[1] = p1;
 
1516
                    tab[2] = fma(p2 - p1, 1-h, p1);
 
1517
                    tab[3] = fma(p2 - p1, h, p1);
 
1518
 
 
1519
                    b = tab[sector_data[sector][0]];
 
1520
                    g = tab[sector_data[sector][1]];
 
1521
                    r = tab[sector_data[sector][2]];
 
1522
                }
 
1523
                else
 
1524
                    b = g = r = l;
 
1525
 
 
1526
                dst[bidx] = b;
 
1527
                dst[1] = g;
 
1528
                dst[bidx^2] = r;
 
1529
#if dcn == 4
 
1530
                dst[3] = MAX_NUM;
 
1531
#endif
 
1532
 
 
1533
                ++y;
 
1534
                dst_index += dst_step;
 
1535
                src_index += src_step;
 
1536
            }
 
1537
        }
 
1538
    }
 
1539
}
 
1540
 
 
1541
#endif
 
1542
 
 
1543
/////////////////////////// RGBA <-> mRGBA (alpha premultiplied) //////////////
 
1544
 
 
1545
#ifdef DEPTH_0
 
1546
 
 
1547
__kernel void RGBA2mRGBA(__global const uchar* src, int src_step, int src_offset,
 
1548
                         __global uchar* dst, int dst_step, int dst_offset,
 
1549
                         int rows, int cols)
 
1550
{
 
1551
    int x = get_global_id(0);
 
1552
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1553
 
 
1554
    if (x < cols)
 
1555
    {
 
1556
        int src_index = mad24(y, src_step, src_offset + (x << 2));
 
1557
        int dst_index = mad24(y, dst_step, dst_offset + (x << 2));
 
1558
 
 
1559
        #pragma unroll
 
1560
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1561
        {
 
1562
            if (y < rows)
 
1563
            {
 
1564
                uchar4 src_pix = *(__global const uchar4 *)(src + src_index);
 
1565
 
 
1566
                *(__global uchar4 *)(dst + dst_index) =
 
1567
                    (uchar4)(mad24(src_pix.x, src_pix.w, HALF_MAX) / MAX_NUM,
 
1568
                             mad24(src_pix.y, src_pix.w, HALF_MAX) / MAX_NUM,
 
1569
                             mad24(src_pix.z, src_pix.w, HALF_MAX) / MAX_NUM, src_pix.w);
 
1570
 
 
1571
                ++y;
 
1572
                dst_index += dst_step;
 
1573
                src_index += src_step;
 
1574
            }
 
1575
        }
 
1576
    }
 
1577
}
 
1578
 
 
1579
__kernel void mRGBA2RGBA(__global const uchar* src, int src_step, int src_offset,
 
1580
                         __global uchar* dst, int dst_step, int dst_offset,
 
1581
                         int rows, int cols)
 
1582
{
 
1583
    int x = get_global_id(0);
 
1584
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1585
 
 
1586
    if (x < cols)
 
1587
    {
 
1588
        int src_index = mad24(y, src_step, mad24(x, 4, src_offset));
 
1589
        int dst_index = mad24(y, dst_step, mad24(x, 4, dst_offset));
 
1590
 
 
1591
        #pragma unroll
 
1592
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1593
        {
 
1594
            if (y < rows)
 
1595
            {
 
1596
                uchar4 src_pix = *(__global const uchar4 *)(src + src_index);
 
1597
                uchar v3 = src_pix.w, v3_half = v3 / 2;
 
1598
 
 
1599
                if (v3 == 0)
 
1600
                    *(__global uchar4 *)(dst + dst_index) = (uchar4)(0, 0, 0, 0);
 
1601
                else
 
1602
                    *(__global uchar4 *)(dst + dst_index) =
 
1603
                        (uchar4)(mad24(src_pix.x, MAX_NUM, v3_half) / v3,
 
1604
                                 mad24(src_pix.y, MAX_NUM, v3_half) / v3,
 
1605
                                 mad24(src_pix.z, MAX_NUM, v3_half) / v3, v3);
 
1606
 
 
1607
                ++y;
 
1608
                dst_index += dst_step;
 
1609
                src_index += src_step;
 
1610
            }
 
1611
        }
 
1612
    }
 
1613
}
 
1614
 
 
1615
#endif
 
1616
 
 
1617
/////////////////////////////////// [l|s]RGB <-> Lab ///////////////////////////
 
1618
 
 
1619
#define lab_shift xyz_shift
 
1620
#define gamma_shift 3
 
1621
#define lab_shift2 (lab_shift + gamma_shift)
 
1622
#define GAMMA_TAB_SIZE 1024
 
1623
#define GammaTabScale (float)GAMMA_TAB_SIZE
 
1624
 
 
1625
inline float splineInterpolate(float x, __global const float * tab, int n)
 
1626
{
 
1627
    int ix = clamp(convert_int_sat_rtn(x), 0, n-1);
 
1628
    x -= ix;
 
1629
    tab += ix << 2;
 
1630
    return fma(fma(fma(tab[3], x, tab[2]), x, tab[1]), x, tab[0]);
 
1631
}
 
1632
 
 
1633
#ifdef DEPTH_0
 
1634
 
 
1635
__kernel void BGR2Lab(__global const uchar * src, int src_step, int src_offset,
 
1636
                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
 
1637
                      __global const ushort * gammaTab, __global ushort * LabCbrtTab_b,
 
1638
                      __constant int * coeffs, int Lscale, int Lshift)
 
1639
{
 
1640
    int x = get_global_id(0);
 
1641
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1642
 
 
1643
    if (x < cols)
 
1644
    {
 
1645
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1646
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1647
 
 
1648
        #pragma unroll
 
1649
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1650
        {
 
1651
            if (y < rows)
 
1652
            {
 
1653
                __global const uchar* src_ptr = src + src_index;
 
1654
                __global uchar* dst_ptr = dst + dst_index;
 
1655
                uchar4 src_pix = vload4(0, src_ptr);
 
1656
 
 
1657
                int C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
 
1658
                    C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
 
1659
                    C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
 
1660
 
 
1661
                int R = gammaTab[src_pix.x], G = gammaTab[src_pix.y], B = gammaTab[src_pix.z];
 
1662
                int fX = LabCbrtTab_b[CV_DESCALE(mad24(R, C0, mad24(G, C1, B*C2)), lab_shift)];
 
1663
                int fY = LabCbrtTab_b[CV_DESCALE(mad24(R, C3, mad24(G, C4, B*C5)), lab_shift)];
 
1664
                int fZ = LabCbrtTab_b[CV_DESCALE(mad24(R, C6, mad24(G, C7, B*C8)), lab_shift)];
 
1665
 
 
1666
                int L = CV_DESCALE( Lscale*fY + Lshift, lab_shift2 );
 
1667
                int a = CV_DESCALE( mad24(500, fX - fY, 128*(1 << lab_shift2)), lab_shift2 );
 
1668
                int b = CV_DESCALE( mad24(200, fY - fZ, 128*(1 << lab_shift2)), lab_shift2 );
 
1669
 
 
1670
                dst_ptr[0] = SAT_CAST(L);
 
1671
                dst_ptr[1] = SAT_CAST(a);
 
1672
                dst_ptr[2] = SAT_CAST(b);
 
1673
 
 
1674
                ++y;
 
1675
                dst_index += dst_step;
 
1676
                src_index += src_step;
 
1677
            }
 
1678
        }
 
1679
    }
 
1680
}
 
1681
 
 
1682
#elif defined DEPTH_5
 
1683
 
 
1684
__kernel void BGR2Lab(__global const uchar * srcptr, int src_step, int src_offset,
 
1685
                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
 
1686
#ifdef SRGB
 
1687
                      __global const float * gammaTab,
 
1688
#endif
 
1689
                      __constant float * coeffs, float _1_3, float _a)
 
1690
{
 
1691
    int x = get_global_id(0);
 
1692
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1693
 
 
1694
    if (x < cols)
 
1695
    {
 
1696
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1697
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1698
 
 
1699
        #pragma unroll
 
1700
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1701
        {
 
1702
            if (y < rows)
 
1703
            {
 
1704
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1705
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1706
                float4 src_pix = vload4(0, src);
 
1707
 
 
1708
                float C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
 
1709
                      C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
 
1710
                      C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
 
1711
 
 
1712
                float R = clamp(src_pix.x, 0.0f, 1.0f);
 
1713
                float G = clamp(src_pix.y, 0.0f, 1.0f);
 
1714
                float B = clamp(src_pix.z, 0.0f, 1.0f);
 
1715
 
 
1716
#ifdef SRGB
 
1717
                R = splineInterpolate(R * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1718
                G = splineInterpolate(G * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1719
                B = splineInterpolate(B * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1720
#endif
 
1721
 
 
1722
                float X = fma(R, C0, fma(G, C1, B*C2));
 
1723
                float Y = fma(R, C3, fma(G, C4, B*C5));
 
1724
                float Z = fma(R, C6, fma(G, C7, B*C8));
 
1725
 
 
1726
                float FX = X > 0.008856f ? rootn(X, 3) : fma(7.787f, X, _a);
 
1727
                float FY = Y > 0.008856f ? rootn(Y, 3) : fma(7.787f, Y, _a);
 
1728
                float FZ = Z > 0.008856f ? rootn(Z, 3) : fma(7.787f, Z, _a);
 
1729
 
 
1730
                float L = Y > 0.008856f ? fma(116.f, FY, -16.f) : (903.3f * Y);
 
1731
                float a = 500.f * (FX - FY);
 
1732
                float b = 200.f * (FY - FZ);
 
1733
 
 
1734
                dst[0] = L;
 
1735
                dst[1] = a;
 
1736
                dst[2] = b;
 
1737
 
 
1738
                ++y;
 
1739
                dst_index += dst_step;
 
1740
                src_index += src_step;
 
1741
            }
 
1742
        }
 
1743
    }
 
1744
}
 
1745
 
 
1746
#endif
 
1747
 
 
1748
inline void Lab2BGR_f(const float * srcbuf, float * dstbuf,
 
1749
#ifdef SRGB
 
1750
                      __global const float * gammaTab,
 
1751
#endif
 
1752
                      __constant float * coeffs, float lThresh, float fThresh)
 
1753
{
 
1754
    float li = srcbuf[0], ai = srcbuf[1], bi = srcbuf[2];
 
1755
 
 
1756
    float C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
 
1757
          C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
 
1758
          C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
 
1759
 
 
1760
    float y, fy;
 
1761
    if (li <= lThresh)
 
1762
    {
 
1763
        y = li / 903.3f;
 
1764
        fy = fma(7.787f, y, 16.0f / 116.0f);
 
1765
    }
 
1766
    else
 
1767
    {
 
1768
        fy = (li + 16.0f) / 116.0f;
 
1769
        y = fy * fy * fy;
 
1770
    }
 
1771
 
 
1772
    float fxz[] = { ai / 500.0f + fy, fy - bi / 200.0f };
 
1773
 
 
1774
    #pragma unroll
 
1775
    for (int j = 0; j < 2; j++)
 
1776
        if (fxz[j] <= fThresh)
 
1777
            fxz[j] = (fxz[j] - 16.0f / 116.0f) / 7.787f;
 
1778
        else
 
1779
            fxz[j] = fxz[j] * fxz[j] * fxz[j];
 
1780
 
 
1781
    float x = fxz[0], z = fxz[1];
 
1782
    float ro = clamp(fma(C0, x, fma(C1, y, C2 * z)), 0.0f, 1.0f);
 
1783
    float go = clamp(fma(C3, x, fma(C4, y, C5 * z)), 0.0f, 1.0f);
 
1784
    float bo = clamp(fma(C6, x, fma(C7, y, C8 * z)), 0.0f, 1.0f);
 
1785
 
 
1786
#ifdef SRGB
 
1787
    ro = splineInterpolate(ro * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1788
    go = splineInterpolate(go * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1789
    bo = splineInterpolate(bo * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1790
#endif
 
1791
 
 
1792
    dstbuf[0] = ro, dstbuf[1] = go, dstbuf[2] = bo;
 
1793
}
 
1794
 
 
1795
#ifdef DEPTH_0
 
1796
 
 
1797
__kernel void Lab2BGR(__global const uchar * src, int src_step, int src_offset,
 
1798
                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
 
1799
#ifdef SRGB
 
1800
                      __global const float * gammaTab,
 
1801
#endif
 
1802
                      __constant float * coeffs, float lThresh, float fThresh)
 
1803
{
 
1804
    int x = get_global_id(0);
 
1805
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1806
 
 
1807
    if (x < cols)
 
1808
    {
 
1809
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1810
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1811
 
 
1812
        #pragma unroll
 
1813
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1814
        {
 
1815
            if (y < rows)
 
1816
            {
 
1817
                __global const uchar* src_ptr = src + src_index;
 
1818
                __global uchar * dst_ptr = dst + dst_index;
 
1819
                uchar4 src_pix = vload4(0, src_ptr);
 
1820
 
 
1821
                float srcbuf[3], dstbuf[3];
 
1822
                srcbuf[0] = src_pix.x*(100.f/255.f);
 
1823
                srcbuf[1] = convert_float(src_pix.y - 128);
 
1824
                srcbuf[2] = convert_float(src_pix.z - 128);
 
1825
 
 
1826
                Lab2BGR_f(&srcbuf[0], &dstbuf[0],
 
1827
#ifdef SRGB
 
1828
                    gammaTab,
 
1829
#endif
 
1830
                    coeffs, lThresh, fThresh);
 
1831
 
 
1832
#if dcn == 3
 
1833
                dst_ptr[0] = SAT_CAST(dstbuf[0] * 255.0f);
 
1834
                dst_ptr[1] = SAT_CAST(dstbuf[1] * 255.0f);
 
1835
                dst_ptr[2] = SAT_CAST(dstbuf[2] * 255.0f);
 
1836
#else
 
1837
                *(__global uchar4 *)dst_ptr = (uchar4)(SAT_CAST(dstbuf[0] * 255.0f),
 
1838
                    SAT_CAST(dstbuf[1] * 255.0f), SAT_CAST(dstbuf[2] * 255.0f), MAX_NUM);
 
1839
#endif
 
1840
                ++y;
 
1841
                dst_index += dst_step;
 
1842
                src_index += src_step;
 
1843
            }
 
1844
        }
 
1845
    }
 
1846
}
 
1847
 
 
1848
#elif defined DEPTH_5
 
1849
 
 
1850
__kernel void Lab2BGR(__global const uchar * srcptr, int src_step, int src_offset,
 
1851
                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
 
1852
#ifdef SRGB
 
1853
                      __global const float * gammaTab,
 
1854
#endif
 
1855
                      __constant float * coeffs, float lThresh, float fThresh)
 
1856
{
 
1857
    int x = get_global_id(0);
 
1858
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1859
 
 
1860
    if (x < cols)
 
1861
    {
 
1862
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1863
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1864
 
 
1865
        #pragma unroll
 
1866
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1867
        {
 
1868
            if (y < rows)
 
1869
            {
 
1870
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1871
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1872
                float4 src_pix = vload4(0, src);
 
1873
 
 
1874
                float srcbuf[3], dstbuf[3];
 
1875
                srcbuf[0] = src_pix.x, srcbuf[1] = src_pix.y, srcbuf[2] = src_pix.z;
 
1876
 
 
1877
                Lab2BGR_f(&srcbuf[0], &dstbuf[0],
 
1878
#ifdef SRGB
 
1879
                    gammaTab,
 
1880
#endif
 
1881
                    coeffs, lThresh, fThresh);
 
1882
 
 
1883
                dst[0] = dstbuf[0], dst[1] = dstbuf[1], dst[2] = dstbuf[2];
 
1884
#if dcn == 4
 
1885
                dst[3] = MAX_NUM;
 
1886
#endif
 
1887
                ++y;
 
1888
                dst_index += dst_step;
 
1889
                src_index += src_step;
 
1890
            }
 
1891
        }
 
1892
    }
 
1893
}
 
1894
 
 
1895
#endif
 
1896
 
 
1897
/////////////////////////////////// [l|s]RGB <-> Luv ///////////////////////////
 
1898
 
 
1899
#define LAB_CBRT_TAB_SIZE 1024
 
1900
#define LAB_CBRT_TAB_SIZE_B (256*3/2*(1<<gamma_shift))
 
1901
 
 
1902
__constant float LabCbrtTabScale = LAB_CBRT_TAB_SIZE/1.5f;
 
1903
 
 
1904
#ifdef DEPTH_5
 
1905
 
 
1906
__kernel void BGR2Luv(__global const uchar * srcptr, int src_step, int src_offset,
 
1907
                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
 
1908
#ifdef SRGB
 
1909
                      __global const float * gammaTab,
 
1910
#endif
 
1911
                      __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
 
1912
{
 
1913
    int x = get_global_id(0);
 
1914
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1915
 
 
1916
    if (x < cols)
 
1917
    {
 
1918
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1919
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1920
 
 
1921
        #pragma unroll
 
1922
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1923
            if (y < rows)
 
1924
            {
 
1925
                __global const float * src = (__global const float *)(srcptr + src_index);
 
1926
                __global float * dst = (__global float *)(dstptr + dst_index);
 
1927
 
 
1928
                float R = src[0], G = src[1], B = src[2];
 
1929
 
 
1930
#ifdef SRGB
 
1931
                R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1932
                G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1933
                B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1934
#endif
 
1935
                float X = fma(R, coeffs[0], fma(G, coeffs[1], B*coeffs[2]));
 
1936
                float Y = fma(R, coeffs[3], fma(G, coeffs[4], B*coeffs[5]));
 
1937
                float Z = fma(R, coeffs[6], fma(G, coeffs[7], B*coeffs[8]));
 
1938
 
 
1939
                float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
 
1940
                L = fma(116.f, L, -16.f);
 
1941
 
 
1942
                float d = 52.0f / fmax(fma(15.0f, Y, fma(3.0f, Z, X)), FLT_EPSILON);
 
1943
                float u = L*fma(X, d, -_un);
 
1944
                float v = L*fma(2.25f, Y*d, -_vn);
 
1945
 
 
1946
                dst[0] = L;
 
1947
                dst[1] = u;
 
1948
                dst[2] = v;
 
1949
 
 
1950
                ++y;
 
1951
                dst_index += dst_step;
 
1952
                src_index += src_step;
 
1953
            }
 
1954
    }
 
1955
}
 
1956
 
 
1957
#elif defined DEPTH_0
 
1958
 
 
1959
__kernel void BGR2Luv(__global const uchar * src, int src_step, int src_offset,
 
1960
                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
 
1961
#ifdef SRGB
 
1962
                      __global const float * gammaTab,
 
1963
#endif
 
1964
                      __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
 
1965
{
 
1966
    int x = get_global_id(0);
 
1967
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
1968
 
 
1969
    if (x < cols)
 
1970
    {
 
1971
        src += mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
1972
        dst += mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
1973
 
 
1974
        #pragma unroll
 
1975
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
1976
            if (y < rows)
 
1977
            {
 
1978
                float scale = 1.0f / 255.0f;
 
1979
                float R = src[0]*scale, G = src[1]*scale, B = src[2]*scale;
 
1980
 
 
1981
#ifdef SRGB
 
1982
                R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1983
                G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1984
                B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
1985
#endif
 
1986
                float X = fma(R, coeffs[0], fma(G, coeffs[1], B*coeffs[2]));
 
1987
                float Y = fma(R, coeffs[3], fma(G, coeffs[4], B*coeffs[5]));
 
1988
                float Z = fma(R, coeffs[6], fma(G, coeffs[7], B*coeffs[8]));
 
1989
 
 
1990
                float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
 
1991
                L = 116.f*L - 16.f;
 
1992
 
 
1993
                float d = (4*13) / fmax(fma(15.0f, Y, fma(3.0f, Z, X)), FLT_EPSILON);
 
1994
                float u = L*(X*d - _un);
 
1995
                float v = L*fma(2.25f, Y*d, -_vn);
 
1996
 
 
1997
                dst[0] = SAT_CAST(L * 2.55f);
 
1998
                dst[1] = SAT_CAST(fma(u, 0.72033898305084743f, 96.525423728813564f));
 
1999
                dst[2] = SAT_CAST(fma(v, 0.9732824427480916f, 136.259541984732824f));
 
2000
 
 
2001
                ++y;
 
2002
                dst += dst_step;
 
2003
                src += src_step;
 
2004
            }
 
2005
    }
 
2006
}
 
2007
 
 
2008
#endif
 
2009
 
 
2010
#ifdef DEPTH_5
 
2011
 
 
2012
__kernel void Luv2BGR(__global const uchar * srcptr, int src_step, int src_offset,
 
2013
                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
 
2014
#ifdef SRGB
 
2015
                      __global const float * gammaTab,
 
2016
#endif
 
2017
                      __constant float * coeffs, float _un, float _vn)
 
2018
{
 
2019
    int x = get_global_id(0);
 
2020
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
2021
 
 
2022
    if (x < cols)
 
2023
    {
 
2024
        int src_index = mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
2025
        int dst_index = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
2026
 
 
2027
        #pragma unroll
 
2028
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
2029
            if (y < rows)
 
2030
            {
 
2031
                __global const float * src = (__global const float *)(srcptr + src_index);
 
2032
                __global float * dst = (__global float *)(dstptr + dst_index);
 
2033
 
 
2034
                float L = src[0], u = src[1], v = src[2], d, X, Y, Z;
 
2035
                Y = (L + 16.f) * (1.f/116.f);
 
2036
                Y = Y*Y*Y;
 
2037
                d = (1.f/13.f)/L;
 
2038
                u = fma(u, d, _un);
 
2039
                v = fma(v, d, _vn);
 
2040
                float iv = 1.f/v;
 
2041
                X = 2.25f * u * Y * iv;
 
2042
                Z = (12 - fma(3.0f, u, 20.0f * v)) * Y * 0.25f * iv;
 
2043
 
 
2044
                float R = fma(X, coeffs[0], fma(Y, coeffs[1], Z * coeffs[2]));
 
2045
                float G = fma(X, coeffs[3], fma(Y, coeffs[4], Z * coeffs[5]));
 
2046
                float B = fma(X, coeffs[6], fma(Y, coeffs[7], Z * coeffs[8]));
 
2047
 
 
2048
                R = clamp(R, 0.f, 1.f);
 
2049
                G = clamp(G, 0.f, 1.f);
 
2050
                B = clamp(B, 0.f, 1.f);
 
2051
 
 
2052
#ifdef SRGB
 
2053
                R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2054
                G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2055
                B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2056
#endif
 
2057
 
 
2058
                dst[0] = R;
 
2059
                dst[1] = G;
 
2060
                dst[2] = B;
 
2061
#if dcn == 4
 
2062
                dst[3] = MAX_NUM;
 
2063
#endif
 
2064
                ++y;
 
2065
                dst_index += dst_step;
 
2066
                src_index += src_step;
 
2067
            }
 
2068
    }
 
2069
}
 
2070
 
 
2071
#elif defined DEPTH_0
 
2072
 
 
2073
__kernel void Luv2BGR(__global const uchar * src, int src_step, int src_offset,
 
2074
                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
 
2075
#ifdef SRGB
 
2076
                      __global const float * gammaTab,
 
2077
#endif
 
2078
                      __constant float * coeffs, float _un, float _vn)
 
2079
{
 
2080
    int x = get_global_id(0);
 
2081
    int y = get_global_id(1) * PIX_PER_WI_Y;
 
2082
 
 
2083
    if (x < cols)
 
2084
    {
 
2085
        src += mad24(y, src_step, mad24(x, scnbytes, src_offset));
 
2086
        dst += mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
 
2087
 
 
2088
        #pragma unroll
 
2089
        for (int cy = 0; cy < PIX_PER_WI_Y; ++cy)
 
2090
            if (y < rows)
 
2091
            {
 
2092
                float d, X, Y, Z;
 
2093
                float L = src[0]*(100.f/255.f);
 
2094
                float u = fma(convert_float(src[1]), 1.388235294117647f, -134.f);
 
2095
                float v = fma(convert_float(src[2]), 1.027450980392157f, - 140.f);
 
2096
                Y = (L + 16.f) * (1.f/116.f);
 
2097
                Y = Y*Y*Y;
 
2098
                d = (1.f/13.f)/L;
 
2099
                u = fma(u, d, _un);
 
2100
                v = fma(v, d, _vn);
 
2101
                float iv = 1.f/v;
 
2102
                X = 2.25f * u * Y * iv ;
 
2103
                Z = (12 - fma(3.0f, u, 20.0f * v)) * Y * 0.25f * iv;
 
2104
 
 
2105
                float R = fma(X, coeffs[0], fma(Y, coeffs[1], Z * coeffs[2]));
 
2106
                float G = fma(X, coeffs[3], fma(Y, coeffs[4], Z * coeffs[5]));
 
2107
                float B = fma(X, coeffs[6], fma(Y, coeffs[7], Z * coeffs[8]));
 
2108
 
 
2109
                R = clamp(R, 0.f, 1.f);
 
2110
                G = clamp(G, 0.f, 1.f);
 
2111
                B = clamp(B, 0.f, 1.f);
 
2112
 
 
2113
#ifdef SRGB
 
2114
                R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2115
                G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2116
                B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
 
2117
#endif
 
2118
 
 
2119
                uchar dst0 = SAT_CAST(R * 255.0f);
 
2120
                uchar dst1 = SAT_CAST(G * 255.0f);
 
2121
                uchar dst2 = SAT_CAST(B * 255.0f);
 
2122
 
 
2123
#if dcn == 4
 
2124
                *(__global uchar4 *)dst = (uchar4)(dst0, dst1, dst2, MAX_NUM);
 
2125
#else
 
2126
                dst[0] = dst0;
 
2127
                dst[1] = dst1;
 
2128
                dst[2] = dst2;
 
2129
#endif
 
2130
 
 
2131
                ++y;
 
2132
                dst += dst_step;
 
2133
                src += src_step;
 
2134
            }
 
2135
    }
 
2136
}
 
2137
 
 
2138
#endif