Skip to content

O
opencv
Commit e52459a3 authored by Andrey Pavlenko's avatar Andrey Pavlenko Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #1708 from ilya-lavrenov:ocl_threshold

parents 106d5eda 1f7f9c96
Hide whitespace changes
Inline Side-by-side
Showing with 68 additions and 152 deletions
modules/ocl/src/imgproc.cpp
View file @ e52459a3
...@@ -98,80 +98,66 @@ namespace cv ...@@ -98,80 +98,66 @@ namespace cv
///////////////////////////////////////////////////////////////////////////////////// /////////////////////////////////////////////////////////////////////////////////////
// threshold // threshold
typedef void (*gpuThresh_t)(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type); static std::vector<uchar> scalarToVector(const cv::Scalar & sc, int depth, int ocn, int cn)
static void threshold_8u(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
{ {
uchar thresh_uchar = cvFloor(thresh); CV_Assert(ocn == cn || (ocn == 4 && cn == 3));
uchar max_val = cvRound(maxVal);
size_t cols = (dst.cols + (dst.offset % 16) + 15) / 16; static const int sizeMap[] = { sizeof(uchar), sizeof(char), sizeof(ushort),
size_t bSizeX = 16, bSizeY = 16; sizeof(short), sizeof(int), sizeof(float), sizeof(double) };
size_t gSizeX = cols % bSizeX == 0 ? cols : (cols + bSizeX - 1) / bSizeX * bSizeX;
size_t gSizeY = dst.rows;
size_t globalThreads[3] = {gSizeX, gSizeY, 1};
size_t localThreads[3] = {bSizeX, bSizeY, 1};
vector< pair<size_t, const void *> > args; int elemSize1 = sizeMap[depth];
args.push_back( make_pair(sizeof(cl_mem), &src.data)); int bufSize = elemSize1 * ocn;
args.push_back( make_pair(sizeof(cl_mem), &dst.data)); std::vector<uchar> _buf(bufSize);
args.push_back( make_pair(sizeof(cl_int), (void *)&src.offset)); uchar * buf = &_buf[0];
args.push_back( make_pair(sizeof(cl_int), (void *)&src.step)); scalarToRawData(sc, buf, CV_MAKE_TYPE(depth, cn));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.offset)); memset(buf + elemSize1 * cn, 0, (ocn - cn) * elemSize1);
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.rows));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.cols)); return _buf;
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.step));
args.push_back( make_pair(sizeof(cl_uchar), (void *)&thresh_uchar));
args.push_back( make_pair(sizeof(cl_uchar), (void *)&max_val));
args.push_back( make_pair(sizeof(cl_int), (void *)&type));
openCLExecuteKernel(src.clCxt, &imgproc_threshold, "threshold", globalThreads, localThreads, args, src.oclchannels(), src.depth());
} }
static void threshold_32f(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type) static void threshold_runner(const oclMat &src, oclMat &dst, double thresh, double maxVal, int thresholdType)
{ {
float thresh_f = thresh; bool ival = src.depth() < CV_32F;
float max_val = maxVal; std::vector<uchar> thresholdValue = scalarToVector(cv::Scalar::all(ival ? cvFloor(thresh) : thresh), dst.depth(),
int dst_offset = (dst.offset >> 2); dst.oclchannels(), dst.channels());
int dst_step = (dst.step >> 2); std::vector<uchar> maxValue = scalarToVector(cv::Scalar::all(maxVal), dst.depth(), dst.oclchannels(), dst.channels());
int src_offset = (src.offset >> 2);
int src_step = (src.step >> 2); size_t localThreads[3] = { 16, 16, 1 };
size_t globalThreads[3] = { dst.cols, dst.rows, 1 };
size_t cols = (dst.cols + (dst_offset & 3) + 3) / 4;
size_t bSizeX = 16, bSizeY = 16; const char * const thresholdMap[] = { "THRESH_BINARY", "THRESH_BINARY_INV", "THRESH_TRUNC",
size_t gSizeX = cols % bSizeX == 0 ? cols : (cols + bSizeX - 1) / bSizeX * bSizeX; "THRESH_TOZERO", "THRESH_TOZERO_INV" };
size_t gSizeY = dst.rows; const char * const channelMap[] = { "", "", "2", "4", "4" };
size_t globalThreads[3] = {gSizeX, gSizeY, 1}; const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" };
size_t localThreads[3] = {bSizeX, bSizeY, 1}; std::string buildOptions = format("-D T=%s%s -D %s", typeMap[src.depth()], channelMap[src.channels()],
thresholdMap[thresholdType]);
int src_step = src.step / src.elemSize(), src_offset = src.offset / src.elemSize();
int dst_step = dst.step / dst.elemSize(), dst_offset = dst.offset / dst.elemSize();
vector< pair<size_t, const void *> > args; vector< pair<size_t, const void *> > args;
args.push_back( make_pair(sizeof(cl_mem), &src.data)); args.push_back( make_pair(sizeof(cl_mem), (void *)&src.data));
args.push_back( make_pair(sizeof(cl_mem), &dst.data));
args.push_back( make_pair(sizeof(cl_int), (void *)&src_offset)); args.push_back( make_pair(sizeof(cl_int), (void *)&src_offset));
args.push_back( make_pair(sizeof(cl_int), (void *)&src_step)); args.push_back( make_pair(sizeof(cl_int), (void *)&src_step));
args.push_back( make_pair(sizeof(cl_mem), (void *)&dst.data));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst_offset)); args.push_back( make_pair(sizeof(cl_int), (void *)&dst_offset));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst_step));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.rows)); args.push_back( make_pair(sizeof(cl_int), (void *)&dst.rows));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst.cols)); args.push_back( make_pair(sizeof(cl_int), (void *)&dst.cols));
args.push_back( make_pair(sizeof(cl_int), (void *)&dst_step)); args.push_back( make_pair(thresholdValue.size(), (void *)&thresholdValue[0]));
args.push_back( make_pair(sizeof(cl_float), (void *)&thresh_f)); args.push_back( make_pair(maxValue.size(), (void *)&maxValue[0]));
args.push_back( make_pair(sizeof(cl_float), (void *)&max_val));
args.push_back( make_pair(sizeof(cl_int), (void *)&type));
openCLExecuteKernel(src.clCxt, &imgproc_threshold, "threshold", globalThreads, localThreads, args, src.oclchannels(), src.depth());
openCLExecuteKernel(src.clCxt, &imgproc_threshold, "threshold", globalThreads, localThreads, args,
-1, -1, buildOptions.c_str());
} }
// threshold: support 8UC1 and 32FC1 data type and five threshold type double threshold(const oclMat &src, oclMat &dst, double thresh, double maxVal, int thresholdType)
double threshold(const oclMat &src, oclMat &dst, double thresh, double maxVal, int type)
{ {
//TODO: These limitations shall be removed later. CV_Assert(thresholdType == THRESH_BINARY || thresholdType == THRESH_BINARY_INV || thresholdType == THRESH_TRUNC
CV_Assert(src.type() == CV_8UC1 || src.type() == CV_32FC1); || thresholdType == THRESH_TOZERO || thresholdType == THRESH_TOZERO_INV);
CV_Assert(type == THRESH_BINARY || type == THRESH_BINARY_INV || type == THRESH_TRUNC
|| type == THRESH_TOZERO || type == THRESH_TOZERO_INV );
static const gpuThresh_t gpuThresh_callers[2] = {threshold_8u, threshold_32f}; dst.create(src.size(), src.type());
threshold_runner(src, dst, thresh, maxVal, thresholdType);
dst.create( src.size(), src.type() );
gpuThresh_callers[(src.type() == CV_32FC1)](src, dst, thresh, maxVal, type);
return thresh; return thresh;
} }
......
modules/ocl/src/opencl/imgproc_threshold.cl
View file @ e52459a3
...@@ -44,109 +44,37 @@ ...@@ -44,109 +44,37 @@
//M*/ //M*/
#if defined (DOUBLE_SUPPORT) #if defined (DOUBLE_SUPPORT)
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#elif defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable #pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif #endif
#endif
// threshold type: __kernel void threshold(__global const T * restrict src, int src_offset, int src_step,
// enum { THRESH_BINARY=0, THRESH_BINARY_INV=1, THRESH_TRUNC=2, THRESH_TOZERO=3, __global T * dst, int dst_offset, int dst_step,
// THRESH_TOZERO_INV=4, THRESH_MASK=7, THRESH_OTSU=8 }; int rows, int cols, T thresh, T max_val)
__kernel void threshold_C1_D0(__global const uchar * restrict src, __global uchar *dst,
int src_offset, int src_step,
int dst_offset, int dst_rows, int dst_cols, int dst_step,
uchar thresh, uchar max_val, int thresh_type
)
{ {
int gx = get_global_id(0); int gx = get_global_id(0);
const int gy = get_global_id(1); int gy = get_global_id(1);
int offset = (dst_offset & 15);
src_offset -= offset;
int dstart = (gx << 4) - offset; if (gx < cols && gy < rows)
if(dstart < dst_cols && gy < dst_rows)
{ {
uchar16 sdata = vload16(gx, src+src_offset+gy*src_step); int src_index = mad24(gy, src_step, src_offset + gx);
uchar16 ddata; int dst_index = mad24(gy, dst_step, dst_offset + gx);
uchar16 zero = 0;
switch (thresh_type)
{
case 0:
ddata = ((sdata > thresh) ) ? (uchar16)(max_val) : (uchar16)(0);
break;
case 1:
ddata = ((sdata > thresh)) ? zero : (uchar16)(max_val);
break;
case 2:
ddata = ((sdata > thresh)) ? (uchar16)(thresh) : sdata;
break;
case 3:
ddata = ((sdata > thresh)) ? sdata : zero;
break;
case 4:
ddata = ((sdata > thresh)) ? zero : sdata;
break;
default:
ddata = sdata;
}
int16 dpos = (int16)(dstart, dstart+1, dstart+2, dstart+3, dstart+4, dstart+5, dstart+6, dstart+7, dstart+8,
dstart+9, dstart+10, dstart+11, dstart+12, dstart+13, dstart+14, dstart+15);
uchar16 dVal = *(__global uchar16*)(dst+dst_offset+gy*dst_step+dstart);
int16 con = dpos >= 0 && dpos < dst_cols;
ddata = convert_uchar16(con != 0) ? ddata : dVal;
if(dstart < dst_cols)
{
*(__global uchar16*)(dst+dst_offset+gy*dst_step+dstart) = ddata;
}
}
}
T sdata = src[src_index], zero = (T)(0);
__kernel void threshold_C1_D5(__global const float * restrict src, __global float *dst, #ifdef THRESH_BINARY
int src_offset, int src_step, dst[dst_index] = sdata > thresh ? max_val : zero;
int dst_offset, int dst_rows, int dst_cols, int dst_step, #elif defined THRESH_BINARY_INV
float thresh, float max_val, int thresh_type dst[dst_index] = sdata > thresh ? zero : max_val;
) #elif defined THRESH_TRUNC
{ dst[dst_index] = sdata > thresh ? thresh : sdata;
const int gx = get_global_id(0); #elif defined THRESH_TOZERO
const int gy = get_global_id(1); dst[dst_index] = sdata > thresh ? sdata : zero;
#elif defined THRESH_TOZERO_INV
int offset = (dst_offset & 3); dst[dst_index] = sdata > thresh ? zero : sdata;
src_offset -= offset; #endif
int dstart = (gx << 2) - offset;
if(dstart < dst_cols && gy < dst_rows)
{
float4 sdata = vload4(gx, src+src_offset+gy*src_step);
float4 ddata;
float4 zero = 0;
switch (thresh_type)
{
case 0:
ddata = sdata > thresh ? (float4)(max_val) : (float4)(0.f);
break;
case 1:
ddata = sdata > thresh ? zero : (float4)max_val;
break;
case 2:
ddata = sdata > thresh ? (float4)thresh : sdata;
break;
case 3:
ddata = sdata > thresh ? sdata : (float4)(0.f);
break;
case 4:
ddata = sdata > thresh ? (float4)(0.f) : sdata;
break;
default:
ddata = sdata;
}
int4 dpos = (int4)(dstart, dstart+1, dstart+2, dstart+3);
float4 dVal = *(__global float4*)(dst+dst_offset+gy*dst_step+dstart);
int4 con = dpos >= 0 && dpos < dst_cols;
ddata = convert_float4(con) != (float4)(0) ? ddata : dVal;
if(dstart < dst_cols)
{
*(__global float4*)(dst+dst_offset+gy*dst_step+dstart) = ddata;
}
} }
} }
modules/ocl/test/test_imgproc.cpp
View file @ e52459a3
...@@ -502,7 +502,9 @@ INSTANTIATE_TEST_CASE_P(Imgproc, Integral, Combine( ...@@ -502,7 +502,9 @@ INSTANTIATE_TEST_CASE_P(Imgproc, Integral, Combine(
Bool())); Bool()));
INSTANTIATE_TEST_CASE_P(Imgproc, Threshold, Combine( INSTANTIATE_TEST_CASE_P(Imgproc, Threshold, Combine(
Values(CV_8UC1, CV_32FC1), Values(CV_8UC1, CV_8UC2, CV_8UC3, CV_8UC4,
CV_16SC1, CV_16SC2, CV_16SC3, CV_16SC4,
CV_32FC1, CV_32FC2, CV_32FC3, CV_32FC4),
Values(0), Values(0),
Values(ThreshOp(THRESH_BINARY), Values(ThreshOp(THRESH_BINARY),
ThreshOp(THRESH_BINARY_INV), ThreshOp(THRESH_TRUNC), ThreshOp(THRESH_BINARY_INV), ThreshOp(THRESH_TRUNC),
......
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