Chiebot-Mirror
/
opencv
8
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Merge pull request #1847 from ilya-lavrenov:ocl_imgproc_fix

Browse Source
pull/1851/head
Andrey Pavlenko 12 years ago committed by OpenCV Buildbot
parent ca8c129c57 561a7f5782
commit c33cb94d66
2 changed files with 32 additions and 32 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 36
      modules/imgproc/src/opencl/cvtcolor.cl
  2. 28
      modules/imgproc/src/opencl/resize.cl

36
modules/imgproc/src/opencl/cvtcolor.cl
Unescape View File

@ -98,9 +98,9 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset,
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes));
DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes));
#if defined (DEPTH_5) #ifdef DEPTH_5
dst[0] = src[bidx] * 0.114f + src[1] * 0.587f + src[(bidx^2)] * 0.299f; dst[0] = src[bidx] * 0.114f + src[1] * 0.587f + src[(bidx^2)] * 0.299f;
#else #else
dst[0] = (DATA_TYPE)CV_DESCALE((src[bidx] * B2Y + src[1] * G2Y + src[(bidx^2)] * R2Y), yuv_shift); dst[0] = (DATA_TYPE)CV_DESCALE((src[bidx] * B2Y + src[1] * G2Y + src[(bidx^2)] * R2Y), yuv_shift);
@ -117,8 +117,8 @@ __kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset,
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes));
DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes));
DATA_TYPE val = src[0]; DATA_TYPE val = src[0];
dst[0] = dst[1] = dst[2] = val; dst[0] = dst[1] = dst[2] = val;
#if dcn == 4 #if dcn == 4
@ -141,11 +141,11 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int srcstep, int srcoffset,
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes));
DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes));
DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2]; DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2];
#if defined (DEPTH_5) #ifdef DEPTH_5
__constant float * coeffs = c_RGB2YUVCoeffs_f; __constant float * coeffs = c_RGB2YUVCoeffs_f;
const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2]; const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2];
const DATA_TYPE U = (b - Y) * coeffs[3] + HALF_MAX; const DATA_TYPE U = (b - Y) * coeffs[3] + HALF_MAX;
@ -176,11 +176,11 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int srcstep, int srcoffset,
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes));
DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes));
DATA_TYPE Y = src[0], U = src[1], V = src[2]; DATA_TYPE Y = src[0], U = src[1], V = src[2];
#if defined (DEPTH_5) #ifdef DEPTH_5
__constant float * coeffs = c_YUV2RGBCoeffs_f; __constant float * coeffs = c_YUV2RGBCoeffs_f;
const float r = Y + (V - HALF_MAX) * coeffs[3]; const float r = Y + (V - HALF_MAX) * coeffs[3];
const float g = Y + (V - HALF_MAX) * coeffs[2] + (U - HALF_MAX) * coeffs[1]; const float g = Y + (V - HALF_MAX) * coeffs[2] + (U - HALF_MAX) * coeffs[1];
@ -217,10 +217,10 @@ __kernel void YUV2RGBA_NV12(__global const uchar* srcptr, int srcstep, int srcof
if (y < rows / 2 && x < cols / 2 ) if (y < rows / 2 && x < cols / 2 )
{ {
__global const uchar* ysrc = srcptr + mad24(y << 1, srcstep, (x << 1) + srcoffset); __global const uchar* ysrc = (__global const uchar*)(srcptr + mad24(y << 1, srcstep, (x << 1) + srcoffset));
__global const uchar* usrc = srcptr + mad24(rows + y, srcstep, (x << 1) + srcoffset); __global const uchar* usrc = (__global const uchar*)(srcptr + mad24(rows + y, srcstep, (x << 1) + srcoffset));
__global uchar* dst1 = dstptr + mad24(y << 1, dststep, x*(dcn*2) + dstoffset); __global uchar* dst1 = (__global uchar*)(dstptr + mad24(y << 1, dststep, x*(dcn*2) + dstoffset));
__global uchar* dst2 = dstptr + mad24((y << 1) + 1, dststep, x*(dcn*2) + dstoffset); __global uchar* dst2 = (__global uchar*)(dstptr + mad24((y << 1) + 1, dststep, x*(dcn*2) + dstoffset));
int Y1 = ysrc[0]; int Y1 = ysrc[0];
int Y2 = ysrc[1]; int Y2 = ysrc[1];
@ -282,11 +282,11 @@ __kernel void RGB2YCrCb(__global const uchar* srcptr, int srcstep, int srcoffset
if (y < rows && x < cols) if (y < rows && x < cols)
{ {
const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes));
DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes));
DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2]; DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2];
#if defined (DEPTH_5) #ifdef DEPTH_5
__constant float * coeffs = c_RGB2YCrCbCoeffs_f; __constant float * coeffs = c_RGB2YCrCbCoeffs_f;
const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2]; const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2];
const DATA_TYPE Cr = (r - Y) * coeffs[3] + HALF_MAX; const DATA_TYPE Cr = (r - Y) * coeffs[3] + HALF_MAX;

28
modules/imgproc/src/opencl/resize.cl
Unescape View File

@ -87,7 +87,7 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
int y_ = INC(y,srcrows); int y_ = INC(y,srcrows);
int x_ = INC(x,srccols); int x_ = INC(x,srccols);
const PIXTYPE* src = (const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)); __global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE));
#if depth == 0 #if depth == 0
u = u * INTER_RESIZE_COEF_SCALE; u = u * INTER_RESIZE_COEF_SCALE;
@ -98,10 +98,10 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
int U1 = rint(INTER_RESIZE_COEF_SCALE - u); int U1 = rint(INTER_RESIZE_COEF_SCALE - u);
int V1 = rint(INTER_RESIZE_COEF_SCALE - v); int V1 = rint(INTER_RESIZE_COEF_SCALE - v);
WORKTYPE data0 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); WORKTYPE data0 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)));
WORKTYPE data1 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE)));
WORKTYPE data2 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE)));
WORKTYPE data3 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE)));
WORKTYPE val = mul24((WORKTYPE)mul24(U1, V1), data0) + mul24((WORKTYPE)mul24(U, V1), data1) + WORKTYPE val = mul24((WORKTYPE)mul24(U1, V1), data0) + mul24((WORKTYPE)mul24(U, V1), data1) +
mul24((WORKTYPE)mul24(U1, V), data2) + mul24((WORKTYPE)mul24(U, V), data3); mul24((WORKTYPE)mul24(U1, V), data2) + mul24((WORKTYPE)mul24(U, V), data3);
@ -109,16 +109,16 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset,
#else #else
float u1 = 1.f-u; float u1 = 1.f-u;
float v1 = 1.f-v; float v1 = 1.f-v;
WORKTYPE data0 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); WORKTYPE data0 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)));
WORKTYPE data1 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE)));
WORKTYPE data2 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE)));
WORKTYPE data3 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE)));
PIXTYPE uval = u1 * v1 * s_data1 + u * v1 * s_data2 + u1 * v *s_data3 + u * v *s_data4; PIXTYPE uval = u1 * v1 * s_data1 + u * v1 * s_data2 + u1 * v *s_data3 + u * v *s_data4;
#endif #endif
if(dx < dstcols && dy < dstrows) if(dx < dstcols && dy < dstrows)
{ {
PIXTYPE* dst = (PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE)); __global PIXTYPE* dst = (__global PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE));
dst[0] = uval; dst[0] = uval;
} }
} }
@ -140,10 +140,10 @@ __kernel void resizeNN(__global const uchar* srcptr, int srcstep, int srcoffset,
F s2 = dy*ify; F s2 = dy*ify;
int sx = min(convert_int_rtz(s1), srccols-1); int sx = min(convert_int_rtz(s1), srccols-1);
int sy = min(convert_int_rtz(s2), srcrows-1); int sy = min(convert_int_rtz(s2), srcrows-1);
PIXTYPE* dst = (PIXTYPE*)(dstptr +
mad24(dy, dststep, dstoffset + dx*PIXSIZE)); __global PIXTYPE* dst = (__global PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE));
const PIXTYPE* src = (const PIXTYPE*)(srcptr + __global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(sy, srcstep, srcoffset + sx*PIXSIZE));
mad24(sy, srcstep, srcoffset + sx*PIXSIZE));
dst[0] = src[0]; dst[0] = src[0];
} }
} }

Write Preview
Loading…
Cancel
Save