generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

ocl::cvtColor support YUV and YCbCr formats

Browse Source
This commit is contained in:
yao
2013-01-24 14:33:28 +08:00
parent 0773ab4d07
commit 4f778436b5
3 changed files with 614 additions and 66 deletions
Download Patch File Download Diff File Expand all files Collapse all files

219
modules/ocl/src/kernels/cvt_color.cl
View File

@@ -16,6 +16,7 @@
//
// @Authors
// Jia Haipeng, jiahaipeng95@gmail.com
// Peng Xiao, pengxiao@multicorewareinc.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
@@ -48,13 +49,33 @@
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#define DATA_TYPE UNDEFINED
#if defined (DEPTH_0)
#undef DATA_TYPE
#define DATA_TYPE uchar
#define MAX_NUM 255
#define HALF_MAX 128
#define SAT_CAST(num) convert_uchar_sat(num)
#endif
#if defined (DEPTH_2)
#undef DATA_TYPE
#define DATA_TYPE ushort
#define MAX_NUM 65535
#define HALF_MAX 32768
#define SAT_CAST(num) convert_ushort_sat(num)
#endif
#if defined (DEPTH_5)
#undef DATA_TYPE
#define DATA_TYPE float
#define MAX_NUM 1.0f
#define HALF_MAX 0.5f
#define SAT_CAST(num) (num)
#endif
#define CV_DESCALE(x,n) (((x) + (1 << ((n)-1))) >> (n))
enum
{
@@ -65,6 +86,7 @@ enum
B2Y = 1868,
BLOCK_SIZE = 256
};
///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
__kernel void RGB2Gray(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
@@ -72,10 +94,203 @@ __kernel void RGB2Gray(int cols,int rows,int src_step,int dst_step,int channels,
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels * sizeof(DATA_TYPE);
int dst_idx = y * dst_step + x * sizeof(DATA_TYPE);
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x;
#if defined (DEPTH_5)
dst[dst_idx] = src[src_idx + bidx] * 0.114f + src[src_idx + 1] * 0.587f + src[src_idx + (bidx^2)] * 0.299f;
#else
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);
#endif
}
}
__kernel void Gray2RGB(int cols,int rows,int src_step,int dst_step,
__global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x;
int dst_idx = y * dst_step + x * 4;
DATA_TYPE val = src[src_idx];
dst[dst_idx++] = val;
dst[dst_idx++] = val;
dst[dst_idx++] = val;
dst[dst_idx] = MAX_NUM;
}
}
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
__constant float c_RGB2YUVCoeffs_f[5] = { 0.114f, 0.587f, 0.299f, 0.492f, 0.877f };
__constant int c_RGB2YUVCoeffs_i[5] = { B2Y, G2Y, R2Y, 8061, 14369 };
__kernel void RGB2YUV(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YUVCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2];
const DATA_TYPE Cr = (rgb[bidx] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx^2] - Y) * coeffs[4] + HALF_MAX;
#else
__constant int * coeffs = c_RGB2YUVCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift);
const int Y = CV_DESCALE(rgb[0] * coeffs[bidx] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx^2], yuv_shift);
const int Cr = CV_DESCALE((rgb[bidx] - Y) * coeffs[3] + delta, yuv_shift);
const int Cb = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[4] + delta, yuv_shift);
#endif
dst[0] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb );
}
}
__constant float c_YUV2RGBCoeffs_f[5] = { 2.032f, -0.395f, -0.581f, 1.140f };
__constant int c_YUV2RGBCoeffs_i[5] = { 33292, -6472, -9519, 18678 };
__kernel void YUV2RGB(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE yuv[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_YUV2RGBCoeffs_f;
const float b = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[3];
const float g = yuv[0] + (yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1];
const float r = yuv[0] + (yuv[1] - HALF_MAX) * coeffs[0];
#else
__constant int * coeffs = c_YUV2RGBCoeffs_i;
const int b = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[3], yuv_shift);
const int g = yuv[0] + CV_DESCALE((yuv[2] - HALF_MAX) * coeffs[2] + (yuv[1] - HALF_MAX) * coeffs[1], yuv_shift);
const int r = yuv[0] + CV_DESCALE((yuv[1] - HALF_MAX) * coeffs[0], yuv_shift);
#endif
dst[bidx^2] = SAT_CAST( b );
dst[1] = SAT_CAST( g );
dst[bidx] = SAT_CAST( r );
}
}
__constant int ITUR_BT_601_CY = 1220542;
__constant int ITUR_BT_601_CUB = 2116026;
__constant int ITUR_BT_601_CUG = -409993;
__constant int ITUR_BT_601_CVG = -852492;
__constant int ITUR_BT_601_CVR = 1673527;
__constant int ITUR_BT_601_SHIFT = 20;
__kernel void YUV2RGBA_NV12(int cols,int rows,int src_step,int dst_step,
int bidx, int width, int height, __global const uchar* src, __global uchar* dst)
{
const int x = get_global_id(0); // max_x = width / 2
const int y = get_global_id(1); // max_y = height/ 2
if (y < height / 2 && x < width / 2 )
{
__global const uchar* ysrc = src + (y << 1) * src_step + (x << 1);
__global const uchar* usrc = src + (height + y) * src_step + (x << 1);
__global uchar* dst1 = dst + (y << 1) * dst_step + (x << 3);
__global uchar* dst2 = dst + ((y << 1) + 1) * dst_step + (x << 3);
int Y1 = ysrc[0];
int Y2 = ysrc[1];
int Y3 = ysrc[src_step];
int Y4 = ysrc[src_step + 1];
int U = usrc[0] - 128;
int V = usrc[1] - 128;
int ruv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVR * V;
int guv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVG * V + ITUR_BT_601_CUG * U;
int buv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CUB * U;
Y1 = max(0, Y1 - 16) * ITUR_BT_601_CY;
dst1[2 - bidx] = convert_uchar_sat((Y1 + ruv) >> ITUR_BT_601_SHIFT);
dst1[1] = convert_uchar_sat((Y1 + guv) >> ITUR_BT_601_SHIFT);
dst1[bidx] = convert_uchar_sat((Y1 + buv) >> ITUR_BT_601_SHIFT);
dst1[3] = 255;
Y2 = max(0, Y2 - 16) * ITUR_BT_601_CY;
dst1[6 - bidx] = convert_uchar_sat((Y2 + ruv) >> ITUR_BT_601_SHIFT);
dst1[5] = convert_uchar_sat((Y2 + guv) >> ITUR_BT_601_SHIFT);
dst1[4 + bidx] = convert_uchar_sat((Y2 + buv) >> ITUR_BT_601_SHIFT);
dst1[7] = 255;
Y3 = max(0, Y3 - 16) * ITUR_BT_601_CY;
dst2[2 - bidx] = convert_uchar_sat((Y3 + ruv) >> ITUR_BT_601_SHIFT);
dst2[1] = convert_uchar_sat((Y3 + guv) >> ITUR_BT_601_SHIFT);
dst2[bidx] = convert_uchar_sat((Y3 + buv) >> ITUR_BT_601_SHIFT);
dst2[3] = 255;
Y4 = max(0, Y4 - 16) * ITUR_BT_601_CY;
dst2[6 - bidx] = convert_uchar_sat((Y4 + ruv) >> ITUR_BT_601_SHIFT);
dst2[5] = convert_uchar_sat((Y4 + guv) >> ITUR_BT_601_SHIFT);
dst2[4 + bidx] = convert_uchar_sat((Y4 + buv) >> ITUR_BT_601_SHIFT);
dst2[7] = 255;
}
}
///////////////////////////////////// RGB <-> YUV //////////////////////////////////////
__constant float c_RGB2YCrCbCoeffs_f[5] = {0.299f, 0.587f, 0.114f, 0.713f, 0.564f};
__constant int c_RGB2YCrCbCoeffs_i[5] = {R2Y, G2Y, B2Y, 11682, 9241};
__kernel void RGB2YCrCb(int cols,int rows,int src_step,int dst_step,int channels,
int bidx, __global const DATA_TYPE* src, __global DATA_TYPE* dst)
{
const int x = get_global_id(0);
const int y = get_global_id(1);
src_step /= sizeof(DATA_TYPE);
dst_step /= sizeof(DATA_TYPE);
if (y < rows && x < cols)
{
int src_idx = y * src_step + x * channels;
int dst_idx = y * dst_step + x * channels;
dst += dst_idx;
const DATA_TYPE rgb[] = {src[src_idx], src[src_idx + 1], src[src_idx + 2]};
#if defined (DEPTH_5)
__constant float * coeffs = c_RGB2YCrCbCoeffs_f;
const DATA_TYPE Y = rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx];
const DATA_TYPE Cr = (rgb[bidx^2] - Y) * coeffs[3] + HALF_MAX;
const DATA_TYPE Cb = (rgb[bidx] - Y) * coeffs[4] + HALF_MAX;
#else
__constant int * coeffs = c_RGB2YCrCbCoeffs_i;
const int delta = HALF_MAX * (1 << yuv_shift);
const int Y = CV_DESCALE(rgb[0] * coeffs[bidx^2] + rgb[1] * coeffs[1] + rgb[2] * coeffs[bidx], yuv_shift);
const int Cr = CV_DESCALE((rgb[bidx^2] - Y) * coeffs[3] + delta, yuv_shift);
const int Cb = CV_DESCALE((rgb[bidx] - Y) * coeffs[4] + delta, yuv_shift);
#endif
dst[0] = SAT_CAST( Y );
dst[1] = SAT_CAST( Cr );
dst[2] = SAT_CAST( Cb );
}
}