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

improved performance of cv::ocl::countNonZero

Browse Source
This commit is contained in:
Ilya Lavrenov 2014-03-13 15:35:53 +04:00
parent 8e79de35a8
commit 3ebfe60008
3 changed files with 27 additions and 33 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

43
modules/ocl/src/arithm.cpp
View File

@ -1263,38 +1263,35 @@ void cv::ocl::minMaxLoc(const oclMat &src, double *minVal, double *maxVal,
///////////////////////////// countNonZero /////////////////////////////////// ///////////////////////////// countNonZero ///////////////////////////////////
////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////
static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int groupnum, string kernelName) static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int groupnum, int vlen)
{ {
int ochannels = src.oclchannels(); int vElemSize = vlen * src.elemSize1();
int all_cols = src.step / src.elemSize(); int src_step = src.step / vElemSize, src_offset = src.offset / vElemSize;
int pre_cols = (src.offset % src.step) / src.elemSize(); int src_cols = src.cols / vlen, total = src.size().area() / vlen;
int sec_cols = all_cols - (src.offset % src.step + src.cols * src.elemSize() - 1) / src.elemSize() - 1;
int invalid_cols = pre_cols + sec_cols;
int cols = all_cols - invalid_cols , elemnum = cols * src.rows;;
int offset = src.offset / src.elemSize();
const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" }; const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" };
const char * const channelMap[] = { " ", " ", "2", "4", "4" }; const char * const channelMap[] = { "", "", "2", "4", "4", "", "", "", "8" };
string buildOptions = format("-D srcT=%s%s -D dstT=int%s", typeMap[src.depth()], channelMap[ochannels], string buildOptions = format("-D srcT=%s%s -D dstT=int%s -D convertToDstT=convert_int%s",
channelMap[ochannels]); typeMap[src.depth()], channelMap[vlen],
channelMap[vlen], channelMap[vlen]);
vector<pair<size_t , const void *> > args; vector<pair<size_t , const void *> > args;
args.push_back( make_pair( sizeof(cl_int) , (void *)&cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&invalid_cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&offset));
args.push_back( make_pair( sizeof(cl_int) , (void *)&elemnum));
args.push_back( make_pair( sizeof(cl_int) , (void *)&groupnum));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data)); args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&src_cols ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&total ));
args.push_back( make_pair( sizeof(cl_int) , (void *)&groupnum ));
args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst )); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst ));
size_t globalThreads[3] = { groupnum * 256, 1, 1 }; size_t globalThreads[3] = { groupnum * 256, 1, 1 };
#ifdef ANDROID #ifdef ANDROID
openCLExecuteKernel(src.clCxt, &arithm_nonzero, kernelName, globalThreads, NULL, openCLExecuteKernel(src.clCxt, &arithm_nonzero, "arithm_op_nonzero", globalThreads, NULL,
args, -1, -1, buildOptions.c_str()); args, -1, -1, buildOptions.c_str());
#else #else
size_t localThreads[3] = { 256, 1, 1 }; size_t localThreads[3] = { 256, 1, 1 };
openCLExecuteKernel(src.clCxt, &arithm_nonzero, kernelName, globalThreads, localThreads, openCLExecuteKernel(src.clCxt, &arithm_nonzero, "arithm_op_nonzero", globalThreads, localThreads,
args, -1, -1, buildOptions.c_str()); args, -1, -1, buildOptions.c_str());
#endif #endif
} }
@ -1311,18 +1308,20 @@ int cv::ocl::countNonZero(const oclMat &src)
return -1; return -1;
} }
int vlen = 8, vElemSize = src.elemSize1() * vlen;
while (src.offset % vElemSize != 0 || src.step % vElemSize != 0 || src.cols % vlen != 0)
vlen >>= 1, vElemSize >>= 1;
size_t groupnum = src.clCxt->getDeviceInfo().maxComputeUnits; size_t groupnum = src.clCxt->getDeviceInfo().maxComputeUnits;
CV_Assert(groupnum != 0); CV_Assert(groupnum != 0);
int dbsize = groupnum; int dbsize = groupnum * vlen;
string kernelName = "arithm_op_nonzero";
AutoBuffer<int> _buf(dbsize); AutoBuffer<int> _buf(dbsize);
int *p = (int*)_buf, nonzero = 0; int *p = (int*)_buf, nonzero = 0;
memset(p, 0, dbsize * sizeof(int)); memset(p, 0, dbsize * sizeof(int));
cl_mem dstBuffer = openCLCreateBuffer(clCxt, CL_MEM_WRITE_ONLY, dbsize * sizeof(int)); cl_mem dstBuffer = openCLCreateBuffer(clCxt, CL_MEM_WRITE_ONLY, dbsize * sizeof(int));
arithmetic_countNonZero_run(src, dstBuffer, groupnum, kernelName); arithmetic_countNonZero_run(src, dstBuffer, groupnum, vlen);
openCLReadBuffer(clCxt, dstBuffer, (void *)p, dbsize * sizeof(int)); openCLReadBuffer(clCxt, dstBuffer, (void *)p, dbsize * sizeof(int));
for (int i = 0; i < dbsize; i++) for (int i = 0; i < dbsize; i++)

13
modules/ocl/src/opencl/arithm_nonzero.cl
View File

@ -52,23 +52,18 @@
/**************************************Count NonZero**************************************/ /**************************************Count NonZero**************************************/
__kernel void arithm_op_nonzero(int cols, int invalid_cols, int offset, int elemnum, int groupnum, __kernel void arithm_op_nonzero(__global srcT * src, int src_step, int src_offset, int src_cols,
__global srcT *src, __global dstT *dst) int total, int groupnum, __global dstT * dst)
{ {
int lid = get_local_id(0); int lid = get_local_id(0);
int gid = get_group_id(0); int gid = get_group_id(0);
int id = get_global_id(0); int id = get_global_id(0);
int idx = offset + id + (id / cols) * invalid_cols;
__local dstT localmem_nonzero[128]; __local dstT localmem_nonzero[128];
dstT nonzero = (dstT)(0); dstT nonzero = (dstT)(0);
srcT zero = (srcT)(0), one = (srcT)(1);
for (int grain = groupnum << 8; id < elemnum; id += grain) for (int grain = groupnum << 8; id < total; id += grain)
{ nonzero += convertToDstT(src[mad24(id / src_cols, src_step, id % src_cols + src_offset)] == (srcT)(0)) ? (dstT)(0) : (dstT)(1);
idx = offset + id + (id / cols) * invalid_cols;
nonzero += src[idx] == zero ? zero : one;
}
if (lid > 127) if (lid > 127)
localmem_nonzero[lid - 128] = nonzero; localmem_nonzero[lid - 128] = nonzero;

4
modules/ocl/test/test_arithm.cpp
View File

@ -198,7 +198,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
Size roiSize = randomSize(1, MAX_VALUE); Size roiSize = randomSize(1, MAX_VALUE);
Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0); Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src1, src1_roi, roiSize, src1Border, type, 2, 11); randomSubMat(src1, src1_roi, roiSize, src1Border, type, -11, 11);
Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0); Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
randomSubMat(src2, src2_roi, roiSize, src2Border, type, -1540, 1740); randomSubMat(src2, src2_roi, roiSize, src2Border, type, -1540, 1740);
@ -1163,7 +1163,7 @@ OCL_TEST_P(CountNonZero, MAT)
int cpures = cv::countNonZero(src1_roi); int cpures = cv::countNonZero(src1_roi);
int gpures = cv::ocl::countNonZero(gsrc1_roi); int gpures = cv::ocl::countNonZero(gsrc1_roi);
EXPECT_DOUBLE_EQ((double)cpures, (double)gpures); EXPECT_EQ(cpures, gpures);
} }
} }