@@ -115,9 +115,42 @@ enum
115
115
BLOCK_SIZE = 256
116
116
};
117
117
118
+ //constants for conversion from/to RGB and Gray, YUV, YCrCb according to BT.601
119
+ #define B2YF 0.114f
120
+ #define G2YF 0.587f
121
+ #define R2YF 0.299f
122
+ //to YCbCr
123
+ #define YCBF 0.564f
124
+ #define YCRF 0.713f
125
+ #define YCBI 9241
126
+ #define YCRI 11682
127
+ //to YUV
128
+ #define B2UF 0.492f
129
+ #define R2VF 0.877f
130
+ #define B2UI 8061
131
+ #define R2VI 14369
132
+ //from YUV
133
+ #define U2BF 2.032f
134
+ #define U2GF -0.395f
135
+ #define V2GF -0.581f
136
+ #define V2RF 1.140f
137
+ #define U2BI 33292
138
+ #define U2GI -6472
139
+ #define V2GI -9519
140
+ #define V2RI 18678
141
+ //from YCrCb
142
+ #define CR2RF 1.403f
143
+ #define CB2GF -0.344f
144
+ #define CR2GF -0.714f
145
+ #define CB2BF 1.773f
146
+ #define CR2RI 22987
147
+ #define CB2GI -5636
148
+ #define CR2GI -11698
149
+ #define CB2BI 29049
150
+
118
151
///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
119
152
120
- __constant float c_RGB2GrayCoeffs_f [3 ] = { 0.114f , 0.587f , 0.299f };
153
+ __constant float c_RGB2GrayCoeffs_f [3 ] = { B2YF , G2YF , R2YF };
121
154
__constant int c_RGB2GrayCoeffs_i [3 ] = { B2Y , G2Y , R2Y };
122
155
123
156
__kernel void RGB2Gray (int cols , int rows , int src_step , int dst_step ,
@@ -135,7 +168,7 @@ __kernel void RGB2Gray(int cols, int rows, int src_step, int dst_step,
135
168
#ifndef INTEL_DEVICE
136
169
137
170
#ifdef DEPTH_5
138
- dst [dst_idx ] = src [src_idx + bidx ] * 0.114f + src [src_idx + 1 ] * 0.587f + src [src_idx + (bidx ^2 )] * 0.299f ;
171
+ dst [dst_idx ] = src [src_idx + bidx ] * B2YF + src [src_idx + 1 ] * G2YF + src [src_idx + (bidx ^2 )] * R2YF ;
139
172
#else
140
173
dst [dst_idx ] = (DATA_TYPE )CV_DESCALE ((src [src_idx + bidx ] * B2Y + src [src_idx + 1 ] * G2Y + src [src_idx + (bidx ^2 )] * R2Y ), yuv_shift );
141
174
#endif
@@ -221,8 +254,8 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step,
221
254
222
255
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
223
256
224
- __constant float c_RGB2YUVCoeffs_f [5 ] = { 0.114f , 0.587f , 0.299f , 0.492f , 0.877f };
225
- __constant int c_RGB2YUVCoeffs_i [5 ] = { B2Y , G2Y , R2Y , 8061 , 14369 };
257
+ __constant float c_RGB2YUVCoeffs_f [5 ] = { B2YF , G2YF , R2YF , B2UF , R2VF };
258
+ __constant int c_RGB2YUVCoeffs_i [5 ] = { B2Y , G2Y , R2Y , B2UI , R2VI };
226
259
227
260
__kernel void RGB2YUV (int cols , int rows , int src_step , int dst_step ,
228
261
__global const DATA_TYPE * src , __global DATA_TYPE * dst ,
@@ -252,18 +285,18 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
252
285
const DATA_TYPE rgb [] = {src_ptr [0 ], src_ptr [1 ], src_ptr [2 ]};
253
286
254
287
#ifdef DEPTH_5
255
- float Y = rgb [0 ] * coeffs [bidx ^ 2 ] + rgb [1 ] * coeffs [1 ] + rgb [2 ] * coeffs [bidx ];
256
- float U = (rgb [bidx ^ 2 ] - Y ) * coeffs [3 ] + HALF_MAX ;
257
- float V = (rgb [bidx ] - Y ) * coeffs [4 ] + HALF_MAX ;
288
+ float Y = rgb [0 ] * coeffs [bidx ] + rgb [1 ] * coeffs [1 ] + rgb [2 ] * coeffs [bidx ^ 2 ];
289
+ float U = (rgb [bidx ] - Y ) * coeffs [3 ] + HALF_MAX ;
290
+ float V = (rgb [bidx ^ 2 ] - Y ) * coeffs [4 ] + HALF_MAX ;
258
291
#else
259
- int Y = CV_DESCALE (rgb [0 ] * coeffs [bidx ^ 2 ] + rgb [1 ] * coeffs [1 ] + rgb [2 ] * coeffs [bidx ], yuv_shift );
260
- int U = CV_DESCALE ((rgb [bidx ^ 2 ] - Y ) * coeffs [3 ] + delta , yuv_shift );
261
- int V = CV_DESCALE ((rgb [bidx ] - Y ) * coeffs [4 ] + delta , yuv_shift );
292
+ int Y = CV_DESCALE (rgb [0 ] * coeffs [bidx ] + rgb [1 ] * coeffs [1 ] + rgb [2 ] * coeffs [bidx ^ 2 ], yuv_shift );
293
+ int U = CV_DESCALE ((rgb [bidx ] - Y ) * coeffs [3 ] + delta , yuv_shift );
294
+ int V = CV_DESCALE ((rgb [bidx ^ 2 ] - Y ) * coeffs [4 ] + delta , yuv_shift );
262
295
#endif
263
296
264
297
dst_ptr [0 ] = SAT_CAST ( Y );
265
- dst_ptr [1 ] = SAT_CAST ( U );
266
- dst_ptr [2 ] = SAT_CAST ( V );
298
+ dst_ptr [1 ] = SAT_CAST ( V ); //sic! store channels as YVU, not YUV
299
+ dst_ptr [2 ] = SAT_CAST ( U );
267
300
}
268
301
#elif (2 == pixels_per_work_item )
269
302
{
@@ -274,24 +307,24 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
274
307
const float2 c1 = r0 .s15 ;
275
308
const float2 c2 = r0 .s26 ;
276
309
277
- const float2 Y = (bidx == 0 ) ? (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ]) : (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ]);
278
- const float2 U = (bidx == 0 ) ? ((c2 - Y ) * coeffs [3 ] + HALF_MAX ) : ((c0 - Y ) * coeffs [3 ] + HALF_MAX );
279
- const float2 V = (bidx == 0 ) ? ((c0 - Y ) * coeffs [4 ] + HALF_MAX ) : ((c2 - Y ) * coeffs [4 ] + HALF_MAX );
310
+ const float2 Y = (bidx == 0 ) ? (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ]) : (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ]);
311
+ const float2 U = (bidx == 0 ) ? ((c0 - Y ) * coeffs [3 ] + HALF_MAX ) : ((c2 - Y ) * coeffs [3 ] + HALF_MAX );
312
+ const float2 V = (bidx == 0 ) ? ((c2 - Y ) * coeffs [4 ] + HALF_MAX ) : ((c0 - Y ) * coeffs [4 ] + HALF_MAX );
280
313
#else
281
314
const int2 c0 = convert_int2 (r0 .s04 );
282
315
const int2 c1 = convert_int2 (r0 .s15 );
283
316
const int2 c2 = convert_int2 (r0 .s26 );
284
317
285
- const int2 yi = (bidx == 0 ) ? CV_DESCALE (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ], yuv_shift ) : CV_DESCALE (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ], yuv_shift );
286
- const int2 ui = (bidx == 0 ) ? CV_DESCALE ((c2 - yi ) * coeffs [3 ] + delta , yuv_shift ) : CV_DESCALE ((c0 - yi ) * coeffs [3 ] + delta , yuv_shift );
287
- const int2 vi = (bidx == 0 ) ? CV_DESCALE ((c0 - yi ) * coeffs [4 ] + delta , yuv_shift ) : CV_DESCALE ((c2 - yi ) * coeffs [4 ] + delta , yuv_shift );
318
+ const int2 yi = (bidx == 0 ) ? CV_DESCALE (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ], yuv_shift ) : CV_DESCALE (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ], yuv_shift );
319
+ const int2 ui = (bidx == 0 ) ? CV_DESCALE ((c0 - yi ) * coeffs [3 ] + delta , yuv_shift ) : CV_DESCALE ((c2 - yi ) * coeffs [3 ] + delta , yuv_shift );
320
+ const int2 vi = (bidx == 0 ) ? CV_DESCALE ((c2 - yi ) * coeffs [4 ] + delta , yuv_shift ) : CV_DESCALE ((c0 - yi ) * coeffs [4 ] + delta , yuv_shift );
288
321
289
322
const VECTOR2 Y = SAT_CAST2 (yi );
290
323
const VECTOR2 U = SAT_CAST2 (ui );
291
324
const VECTOR2 V = SAT_CAST2 (vi );
292
325
#endif
293
-
294
- vstore8 ((VECTOR8 )(Y .s0 , U .s0 , V .s0 , 0 , Y .s1 , U .s1 , V .s1 , 0 ), 0 , dst_ptr );
326
+ //sic! store channels as YVU, not YUV
327
+ vstore8 ((VECTOR8 )(Y .s0 , V .s0 , U .s0 , 0 , Y .s1 , V .s1 , U .s1 , 0 ), 0 , dst_ptr );
295
328
}
296
329
#elif (4 == pixels_per_work_item )
297
330
{
@@ -302,23 +335,23 @@ __kernel void RGB2YUV(int cols, int rows, int src_step, int dst_step,
302
335
const int4 c1 = convert_int4 (r0 .s159d );
303
336
const int4 c2 = convert_int4 (r0 .s26ae );
304
337
305
- const int4 yi = (bidx == 0 ) ? CV_DESCALE (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ], yuv_shift ) : CV_DESCALE (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ], yuv_shift );
306
- const int4 ui = (bidx == 0 ) ? CV_DESCALE ((c2 - yi ) * coeffs [3 ] + delta , yuv_shift ) : CV_DESCALE ((c0 - yi ) * coeffs [3 ] + delta , yuv_shift );
307
- const int4 vi = (bidx == 0 ) ? CV_DESCALE ((c0 - yi ) * coeffs [4 ] + delta , yuv_shift ) : CV_DESCALE ((c2 - yi ) * coeffs [4 ] + delta , yuv_shift );
338
+ const int4 yi = (bidx == 0 ) ? CV_DESCALE (c0 * coeffs [0 ] + c1 * coeffs [1 ] + c2 * coeffs [2 ], yuv_shift ) : CV_DESCALE (c0 * coeffs [2 ] + c1 * coeffs [1 ] + c2 * coeffs [0 ], yuv_shift );
339
+ const int4 ui = (bidx == 0 ) ? CV_DESCALE ((c0 - yi ) * coeffs [3 ] + delta , yuv_shift ) : CV_DESCALE ((c2 - yi ) * coeffs [3 ] + delta , yuv_shift );
340
+ const int4 vi = (bidx == 0 ) ? CV_DESCALE ((c2 - yi ) * coeffs [4 ] + delta , yuv_shift ) : CV_DESCALE ((c0 - yi ) * coeffs [4 ] + delta , yuv_shift );
308
341
309
342
const VECTOR4 Y = SAT_CAST4 (yi );
310
343
const VECTOR4 U = SAT_CAST4 (ui );
311
344
const VECTOR4 V = SAT_CAST4 (vi );
312
-
313
- vstore16 ((VECTOR16 )(Y .s0 , U .s0 , V .s0 , 0 , Y .s1 , U .s1 , V .s1 , 0 , Y .s2 , U .s2 , V .s2 , 0 , Y .s3 , U .s3 , V .s3 , 0 ), 0 , dst_ptr );
345
+ //sic! store channels as YVU, not YUV
346
+ vstore16 ((VECTOR16 )(Y .s0 , V .s0 , U .s0 , 0 , Y .s1 , V .s1 , U .s1 , 0 , Y .s2 , V .s2 , U .s2 , 0 , Y .s3 , V .s3 , U .s3 , 0 ), 0 , dst_ptr );
314
347
#endif
315
348
}
316
349
#endif //pixels_per_work_item
317
350
}
318
351
}
319
352
320
- __constant float c_YUV2RGBCoeffs_f [5 ] = { 2.032f , -0.395f , -0.581f , 1.140f };
321
- __constant int c_YUV2RGBCoeffs_i [5 ] = { 33292 , -6472 , -9519 , 18678 };
353
+ __constant float c_YUV2RGBCoeffs_f [5 ] = { U2BF , U2GF , V2GF , V2RF };
354
+ __constant int c_YUV2RGBCoeffs_i [5 ] = { U2BI , U2GI , V2GI , V2RI };
322
355
323
356
__kernel void YUV2RGB (int cols , int rows , int src_step , int dst_step ,
324
357
__global const DATA_TYPE * src , __global DATA_TYPE * dst ,
@@ -344,16 +377,17 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
344
377
345
378
#if (1 == pixels_per_work_item )
346
379
{
347
- const DATA_TYPE yuv [] = {src_ptr [0 ], src_ptr [1 ], src_ptr [2 ]};
380
+ //sic! channels stored as YVU, not YUV
381
+ const DATA_TYPE yuv [] = {src_ptr [0 ], src_ptr [2 ], src_ptr [1 ]};
348
382
349
383
#ifdef DEPTH_5
350
- float B = yuv [0 ] + (yuv [2 ] - HALF_MAX ) * coeffs [3 ];
351
- float G = yuv [0 ] + (yuv [2 ] - HALF_MAX ) * coeffs [2 ] + (yuv [1 ] - HALF_MAX ) * coeffs [1 ];
352
- float R = yuv [0 ] + (yuv [1 ] - HALF_MAX ) * coeffs [0 ];
384
+ float B = yuv [0 ] + (yuv [1 ] - HALF_MAX ) * coeffs [0 ];
385
+ float G = yuv [0 ] + (yuv [1 ] - HALF_MAX ) * coeffs [1 ] + (yuv [2 ] - HALF_MAX ) * coeffs [2 ];
386
+ float R = yuv [0 ] + (yuv [2 ] - HALF_MAX ) * coeffs [3 ];
353
387
#else
354
- int B = yuv [0 ] + CV_DESCALE ((yuv [2 ] - HALF_MAX ) * coeffs [3 ], yuv_shift );
355
- int G = yuv [0 ] + CV_DESCALE ((yuv [2 ] - HALF_MAX ) * coeffs [2 ] + (yuv [1 ] - HALF_MAX ) * coeffs [1 ], yuv_shift );
356
- int R = yuv [0 ] + CV_DESCALE ((yuv [1 ] - HALF_MAX ) * coeffs [0 ], yuv_shift );
388
+ int B = yuv [0 ] + CV_DESCALE ((yuv [1 ] - HALF_MAX ) * coeffs [0 ], yuv_shift );
389
+ int G = yuv [0 ] + CV_DESCALE ((yuv [1 ] - HALF_MAX ) * coeffs [1 ] + (yuv [2 ] - HALF_MAX ) * coeffs [2 ], yuv_shift );
390
+ int R = yuv [0 ] + CV_DESCALE ((yuv [2 ] - HALF_MAX ) * coeffs [3 ], yuv_shift );
357
391
#endif
358
392
359
393
dst_ptr [bidx ] = SAT_CAST ( B );
@@ -368,21 +402,23 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
368
402
const VECTOR8 r0 = vload8 (0 , src_ptr );
369
403
370
404
#ifdef DEPTH_5
405
+ //sic! channels stored as YVU, not YUV
371
406
const float2 Y = r0 .s04 ;
372
- const float2 U = r0 .s15 ;
373
- const float2 V = r0 .s26 ;
407
+ const float2 U = r0 .s26 ;
408
+ const float2 V = r0 .s15 ;
374
409
375
- const float2 c0 = (bidx = = 0 ) ? (Y + (V - HALF_MAX ) * coeffs [3 ]) : (Y + (U - HALF_MAX ) * coeffs [0 ]);
410
+ const float2 c0 = (bidx ! = 0 ) ? (Y + (V - HALF_MAX ) * coeffs [3 ]) : (Y + (U - HALF_MAX ) * coeffs [0 ]);
376
411
const float2 c1 = Y + (V - HALF_MAX ) * coeffs [2 ] + (U - HALF_MAX ) * coeffs [1 ];
377
- const float2 c2 = (bidx = = 0 ) ? (Y + (U - HALF_MAX ) * coeffs [0 ]) : (Y + (V - HALF_MAX ) * coeffs [3 ]);
412
+ const float2 c2 = (bidx ! = 0 ) ? (Y + (U - HALF_MAX ) * coeffs [0 ]) : (Y + (V - HALF_MAX ) * coeffs [3 ]);
378
413
#else
414
+ //sic! channels stored as YVU, not YUV
379
415
const int2 Y = convert_int2 (r0 .s04 );
380
- const int2 U = convert_int2 (r0 .s15 );
381
- const int2 V = convert_int2 (r0 .s26 );
416
+ const int2 U = convert_int2 (r0 .s26 );
417
+ const int2 V = convert_int2 (r0 .s15 );
382
418
383
- const int2 c0i = (bidx = = 0 ) ? (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift )) : (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift ));
419
+ const int2 c0i = (bidx ! = 0 ) ? (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift )) : (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift ));
384
420
const int2 c1i = Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [2 ] + (U - HALF_MAX ) * coeffs [1 ], yuv_shift );
385
- const int2 c2i = (bidx = = 0 ) ? (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift )) : (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift ));
421
+ const int2 c2i = (bidx ! = 0 ) ? (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift )) : (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift ));
386
422
387
423
const VECTOR2 c0 = SAT_CAST2 (c0i );
388
424
const VECTOR2 c1 = SAT_CAST2 (c1i );
@@ -400,13 +436,14 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
400
436
#ifndef DEPTH_5
401
437
const VECTOR16 r0 = vload16 (0 , src_ptr );
402
438
439
+ //sic! channels stored as YVU, not YUV
403
440
const int4 Y = convert_int4 (r0 .s048c );
404
- const int4 U = convert_int4 (r0 .s159d );
405
- const int4 V = convert_int4 (r0 .s26ae );
441
+ const int4 U = convert_int4 (r0 .s26ae );
442
+ const int4 V = convert_int4 (r0 .s159d );
406
443
407
- const int4 c0i = (bidx = = 0 ) ? (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift )) : (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift ));
444
+ const int4 c0i = (bidx ! = 0 ) ? (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift )) : (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift ));
408
445
const int4 c1i = Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [2 ] + (U - HALF_MAX ) * coeffs [1 ], yuv_shift );
409
- const int4 c2i = (bidx = = 0 ) ? (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift )) : (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift ));
446
+ const int4 c2i = (bidx ! = 0 ) ? (Y + CV_DESCALE ((U - HALF_MAX ) * coeffs [0 ], yuv_shift )) : (Y + CV_DESCALE ((V - HALF_MAX ) * coeffs [3 ], yuv_shift ));
410
447
411
448
const VECTOR4 c0 = SAT_CAST4 (c0i );
412
449
const VECTOR4 c1 = SAT_CAST4 (c1i );
@@ -484,8 +521,8 @@ __kernel void YUV2RGBA_NV12(int cols, int rows, int src_step, int dst_step,
484
521
485
522
///////////////////////////////////// RGB <-> YCrCb //////////////////////////////////////
486
523
487
- __constant float c_RGB2YCrCbCoeffs_f [5 ] = {0.299f , 0.587f , 0.114f , 0.713f , 0.564f };
488
- __constant int c_RGB2YCrCbCoeffs_i [5 ] = {R2Y , G2Y , B2Y , 11682 , 9241 };
524
+ __constant float c_RGB2YCrCbCoeffs_f [5 ] = {R2YF , G2YF , B2YF , YCRF , YCBF };
525
+ __constant int c_RGB2YCrCbCoeffs_i [5 ] = {R2Y , G2Y , B2Y , YCRI , YCBI };
489
526
490
527
__kernel void RGB2YCrCb (int cols , int rows , int src_step , int dst_step ,
491
528
__global const DATA_TYPE * src , __global DATA_TYPE * dst ,
@@ -579,8 +616,8 @@ __kernel void RGB2YCrCb(int cols, int rows, int src_step, int dst_step,
579
616
}
580
617
}
581
618
582
- __constant float c_YCrCb2RGBCoeffs_f [4 ] = { 1.403f , -0.714f , -0.344f , 1.773f };
583
- __constant int c_YCrCb2RGBCoeffs_i [4 ] = { 22987 , -11698 , -5636 , 29049 };
619
+ __constant float c_YCrCb2RGBCoeffs_f [4 ] = { CR2RF , CR2GF , CB2GF , CB2BF };
620
+ __constant int c_YCrCb2RGBCoeffs_i [4 ] = { CR2RI , CR2GI , CB2GI , CB2BI };
584
621
585
622
__kernel void YCrCb2RGB (int cols , int rows , int src_step , int dst_step ,
586
623
__global const DATA_TYPE * src , __global DATA_TYPE * dst ,
0 commit comments