Skip to content

O
Opencv

Greenplum / Opencv
大约 1 年 前同步成功

通知 7
Star 0
Fork 0
O
Opencv

体验新版 GitCode，发现更多精彩内容 >>

提交 561a7f57 编写于 作者: I Ilya Lavrenov
浏览文件
操作

fixed kernel compilation in imgproc module

上级 9bf2516e
隐藏空白更改
内联 并排
Showing with 32 addition and 32 deletion
modules/imgproc/src/opencl/cvtcolor.cl
浏览文件 @ 561a7f57
...@@ -98,9 +98,9 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, ...@@ -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, ...@@ -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, ...@@ -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, ...@@ -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 ...@@ -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 ...@@ -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;
......
modules/imgproc/src/opencl/resize.cl
浏览文件 @ 561a7f57
...@@ -87,7 +87,7 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset, ...@@ -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, ...@@ -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, ...@@ -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, ...@@ -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];
} }
} }
......
Markdown is supported
0% 或 .
You are about to add 0 people to the discussion. Proceed with caution.
先完成此消息的编辑！
想要评论请 注册 或