Skip to content

O
opencv
Commit 72727111 authored by vbystricky's avatar vbystricky
Browse files
Options

Use 4 pixels for one unit. Some ocl code refactoring

parent 47b092e5
Hide whitespace changes
Inline Side-by-side
Showing with 101 additions and 139 deletions
modules/core/src/convert.cpp
View file @ 72727111
...@@ -1547,30 +1547,21 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst) ...@@ -1547,30 +1547,21 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
int sdepth = _src.depth(); int sdepth = _src.depth();
UMat src = _src.getUMat(), lut = _lut.getUMat(); UMat src = _src.getUMat(), lut = _lut.getUMat();
int dtype = CV_MAKETYPE(ddepth, dcn); _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn));
_dst.create(src.size(), dtype);
UMat dst = _dst.getUMat(); UMat dst = _dst.getUMat();
size_t globalSize[2] = { dst.cols, dst.rows / 2}; size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4};
cv::String build_opt = format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn, ocl::Kernel k("LUT", ocl::core::lut_oclsrc, format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn,
ocl::typeToStr(sdepth), ocl::memopTypeToStr(ddepth) ocl::typeToStr(sdepth), ocl::memopTypeToStr(ddepth)
); ));
if (k.empty())
ocl::Kernel kernel;
if ((4 == lcn) && (CV_8U == sdepth))
kernel.create("LUTC4", ocl::core::lut_oclsrc, build_opt);
else if ((3 == lcn) && (CV_8U == sdepth))
kernel.create("LUTC3", ocl::core::lut_oclsrc, build_opt);
else
kernel.create("LUT", ocl::core::lut_oclsrc, build_opt);
if (kernel.empty())
return false; return false;
kernel.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut), k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
ocl::KernelArg::WriteOnly(dst)); ocl::KernelArg::WriteOnly(dst));
return kernel.run(2, globalSize, NULL, true); return k.run(2, globalSize, NULL, false);
} }
#endif #endif
......
modules/core/src/opencl/lut.cl
View file @ 72727111
...@@ -34,149 +34,120 @@ ...@@ -34,149 +34,120 @@
// //
// //
__kernel void LUTC4(__global const uchar * srcptr, int src_step, int src_offset, #if lcn == 1
__global const uchar * lutptr, int lut_step, int lut_offset, #if dcn == 4
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) #define LUT_OP(num)\
{ uchar4 idx = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
int x = get_global_id(0); dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
int y = 2 * get_global_id(1); dst[0] = lut_l[idx.x];\
dst[1] = lut_l[idx.y];\
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); dst[2] = lut_l[idx.z];\
__local dstT lut_l[256 * lcn];
int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
int step = get_local_size(0) * get_local_size(1);
for (int i = init; i < 256 * lcn; i += step)
{
lut_l[i + 0] = lut[i + 0];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (x < cols && y < rows)
{
int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
__global const uchar4 * src = (__global const uchar4 *)(srcptr + src_index);
int4 idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3);
__global dstT * dst = (__global dstT *)(dstptr + dst_index);
dst[0] = lut_l[idx.x];
dst[1] = lut_l[idx.y];
dst[2] = lut_l[idx.z];
dst[3] = lut_l[idx.w];
if (y < rows - 1)
{
src = (__global const uchar4 *)(srcptr + src_index + src_step);
idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3);
dst = (__global dstT *)(dstptr + dst_index + dst_step);
dst[0] = lut_l[idx.x];
dst[1] = lut_l[idx.y];
dst[2] = lut_l[idx.z];
dst[3] = lut_l[idx.w]; dst[3] = lut_l[idx.w];
} #elif dcn == 3
} #define LUT_OP(num)\
} uchar3 idx = vload3(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
__kernel void LUTC3(__global const uchar * srcptr, int src_step, int src_offset, dst[0] = lut_l[idx.x];\
__global const uchar * lutptr, int lut_step, int lut_offset, dst[1] = lut_l[idx.y];\
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) dst[2] = lut_l[idx.z];
{ #elif dcn == 2
int x = get_global_id(0); #define LUT_OP(num)\
int y = 2 * get_global_id(1); uchar2 idx = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); dst[0] = lut_l[idx.x];\
__local dstT lut_l[256 * lcn];
int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
int step = get_local_size(0) * get_local_size(1);
for (int i = init; i < 256 * lcn; i += step)
{
lut_l[i + 0] = lut[i + 0];
}
barrier(CLK_LOCAL_MEM_FENCE);
if (x < cols && y < rows)
{
int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index));
int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);
__global dstT * dst = (__global dstT *)(dstptr + dst_index);
dst[0] = lut_l[idx.x];
dst[1] = lut_l[idx.y];
dst[2] = lut_l[idx.z];
if (y < rows - 1)
{
uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index + src_step));
idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);
dst = (__global dstT *)(dstptr + dst_index + dst_step);
dst[0] = lut_l[idx.x];
dst[1] = lut_l[idx.y]; dst[1] = lut_l[idx.y];
#elif dcn == 1
#define LUT_OP(num)\
uchar idx = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
dst[0] = lut_l[idx];
#else
#define LUT_OP(num)\
src = (__global const srcT *)(srcptr + src_index + num * src_step);\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
for (int cn = 0; cn < dcn; ++cn)\
dst[cn] = lut_l[src[cn]];
#endif
#else
#if dcn == 4
#define LUT_OP(num)\
uchar4 src_pixel = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
int4 idx = convert_int4(src_pixel) * lcn + (int4)(0, 1, 2, 3);\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
dst[0] = lut_l[idx.x];\
dst[1] = lut_l[idx.y];\
dst[2] = lut_l[idx.z];\
dst[3] = lut_l[idx.w];
#elif dcn == 3
#define LUT_OP(num)\
uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
dst[0] = lut_l[idx.x];\
dst[1] = lut_l[idx.y];\
dst[2] = lut_l[idx.z]; dst[2] = lut_l[idx.z];
} #elif dcn == 2
#define LUT_OP(num)\
uchar2 src_pixel = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\
int2 idx = convert_int2(src_pixel) * lcn + (int2)(0, 1);\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
dst[0] = lut_l[idx.x];\
dst[1] = lut_l[idx.y];
#elif dcn == 1 //error case (1 < lcn) ==> lcn == scn == dcn
#define LUT_OP(num)\
uchar idx = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
dst[0] = lut_l[idx];
#else
#define LUT_OP(num)\
src = (__global const srcT *)(srcptr + src_index + num * src_step);\
dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
for (int cn = 0; cn < dcn; ++cn)\
dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
#endif
#endif
#define LOCAL_LUT_INIT\
{\
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);\
int init = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0));\
int step = get_local_size(0) * get_local_size(1);\
for (int i = init; i < 256 * lcn; i += step)\
{\
lut_l[i] = lut[i];\
}\
barrier(CLK_LOCAL_MEM_FENCE);\
} }
}
__kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset, __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
__global const uchar * lutptr, int lut_step, int lut_offset, __global const uchar * lutptr, int lut_step, int lut_offset,
__global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols) __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols)
{ {
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
__local dstT lut_l[256 * lcn]; __local dstT lut_l[256 * lcn];
int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0)); LOCAL_LUT_INIT;
int step = get_local_size(0) * get_local_size(1);
for (int i = init; i < 256 * lcn; i += step)
{
lut_l[i + 0] = lut[i + 0];
}
barrier(CLK_LOCAL_MEM_FENCE);
int x = get_global_id(0); int x = get_global_id(0);
int y = 2 * get_global_id(1); int y = 4 * get_global_id(1);
if (x < cols && y < rows) if (x < cols && y < rows)
{ {
int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset)); int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
__global const srcT * src = (__global const srcT *)(srcptr + src_index);
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset)); int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
__global dstT * dst = (__global dstT *)(dstptr + dst_index); __global const srcT * src; __global dstT * dst;
#if lcn == 1 LUT_OP(0);
#pragma unroll
for (int cn = 0; cn < dcn; ++cn)
dst[cn] = lut_l[src[cn]];
#else //lcn == scn == dcn
#pragma unroll
for (int cn = 0; cn < dcn; ++cn)
dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
#endif
if (y < rows - 1) if (y < rows - 1)
{ {
src = (__global const srcT *)(srcptr + src_index + src_step); LUT_OP(1);
dst = (__global dstT *)(dstptr + dst_index + dst_step); if (y < rows - 2)
{
#if lcn == 1 LUT_OP(2);
#pragma unroll if (y < rows - 3)
for (int cn = 0; cn < dcn; ++cn) {
dst[cn] = lut_l[src[cn]]; LUT_OP(3);
#else //lcn == scn == dcn }
#pragma unroll }
for (int cn = 0; cn < dcn; ++cn)
dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
#endif
} }
} }
} }
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment