From: Ilya Lavrenov Date: Fri, 27 Sep 2013 13:32:31 +0000 (+0400) Subject: fixed ocl::countNonZero X-Git-Tag: accepted/tizen/6.0/unified/20201030.111113~1314^2~973^2~6 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=b54228fb8332c8537738cd05be2a79cd334b54a2;p=platform%2Fupstream%2Fopencv.git fixed ocl::countNonZero --- diff --git a/modules/ocl/src/arithm.cpp b/modules/ocl/src/arithm.cpp index 0962f92..24420f4 100644 --- a/modules/ocl/src/arithm.cpp +++ b/modules/ocl/src/arithm.cpp @@ -1209,21 +1209,22 @@ void cv::ocl::minMaxLoc(const oclMat &src, double *minVal, double *maxVal, ///////////////////////////// countNonZero /////////////////////////////////// ////////////////////////////////////////////////////////////////////////////// -static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int vlen , int groupnum, string kernelName) +static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int groupnum, string kernelName) { - vector > args; - int all_cols = src.step / (vlen * src.elemSize1()); - int pre_cols = (src.offset % src.step) / (vlen * src.elemSize1()); - int sec_cols = all_cols - (src.offset % src.step + src.cols * src.elemSize() - 1) / (vlen * src.elemSize1()) - 1; + int ochannels = src.oclchannels(); + int all_cols = src.step / src.elemSize(); + int pre_cols = (src.offset % src.step) / src.elemSize(); + 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 / (vlen * src.elemSize1()); - int repeat_s = src.offset / src.elemSize1() - offset * vlen; - int repeat_e = (offset + cols) * vlen - src.offset / src.elemSize1() - src.cols * src.oclchannels(); + int offset = src.offset / src.elemSize(); - char build_options[50]; - sprintf(build_options, "-D DEPTH_%d -D REPEAT_S%d -D REPEAT_E%d", src.depth(), repeat_s, repeat_e); + const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" }; + const char * const channelMap[] = { " ", " ", "2", "4", "4" }; + string buildOptions = format("-D srcT=%s%s -D dstT=int%s", typeMap[src.depth()], channelMap[ochannels], + channelMap[ochannels]); + vector > 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)); @@ -1231,33 +1232,44 @@ static void arithmetic_countNonZero_run(const oclMat &src, cl_mem &dst, int vlen 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 *)&dst )); - size_t gt[3] = {groupnum * 256, 1, 1}, lt[3] = {256, 1, 1}; - openCLExecuteKernel(src.clCxt, &arithm_nonzero, kernelName, gt, lt, args, -1, -1, build_options); + + size_t globalThreads[3] = { groupnum * 256, 1, 1 }; + size_t localThreads[3] = { 256, 1, 1 }; + + openCLExecuteKernel(src.clCxt, &arithm_nonzero, kernelName, globalThreads, localThreads, + args, -1, -1, buildOptions.c_str()); } int cv::ocl::countNonZero(const oclMat &src) { - size_t groupnum = src.clCxt->computeUnits(); + CV_Assert(src.step % src.elemSize() == 0); + CV_Assert(src.channels() == 1); + + Context *clCxt = src.clCxt; if (!src.clCxt->supportsFeature(Context::CL_DOUBLE) && src.depth() == CV_64F) { CV_Error(CV_GpuNotSupported, "selected device doesn't support double"); } + + size_t groupnum = src.clCxt->computeUnits(); CV_Assert(groupnum != 0); - int vlen = 8 , dbsize = groupnum * vlen; - Context *clCxt = src.clCxt; + int dbsize = groupnum; + string kernelName = "arithm_op_nonzero"; AutoBuffer _buf(dbsize); int *p = (int*)_buf, nonzero = 0; - cl_mem dstBuffer = openCLCreateBuffer(clCxt, CL_MEM_WRITE_ONLY, dbsize * sizeof(int)); - arithmetic_countNonZero_run(src, dstBuffer, vlen, groupnum, kernelName); - memset(p, 0, dbsize * sizeof(int)); + + cl_mem dstBuffer = openCLCreateBuffer(clCxt, CL_MEM_WRITE_ONLY, dbsize * sizeof(int)); + arithmetic_countNonZero_run(src, dstBuffer, groupnum, kernelName); openCLReadBuffer(clCxt, dstBuffer, (void *)p, dbsize * sizeof(int)); + for (int i = 0; i < dbsize; i++) nonzero += p[i]; openCLSafeCall(clReleaseMemObject(dstBuffer)); + return nonzero; } diff --git a/modules/ocl/src/opencl/arithm_nonzero.cl b/modules/ocl/src/opencl/arithm_nonzero.cl index e34207d..921367b 100644 --- a/modules/ocl/src/opencl/arithm_nonzero.cl +++ b/modules/ocl/src/opencl/arithm_nonzero.cl @@ -41,151 +41,53 @@ // or tort (including negligence or otherwise) arising in any way out of // the use of this software, even if advised of the possibility of such damage. // -/// -/**************************************PUBLICFUNC*************************************/ #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 #endif - -#if defined (DEPTH_0) -#define VEC_TYPE uchar8 -#endif -#if defined (DEPTH_1) -#define VEC_TYPE char8 -#endif -#if defined (DEPTH_2) -#define VEC_TYPE ushort8 -#endif -#if defined (DEPTH_3) -#define VEC_TYPE short8 -#endif -#if defined (DEPTH_4) -#define VEC_TYPE int8 -#endif -#if defined (DEPTH_5) -#define VEC_TYPE float8 -#endif -#if defined (DEPTH_6) -#define VEC_TYPE double8 #endif -#if defined (REPEAT_S0) -#define repeat_s(a) a = a; -#endif -#if defined (REPEAT_S1) -#define repeat_s(a) a.s0 = 0; -#endif -#if defined (REPEAT_S2) -#define repeat_s(a) a.s0 = 0;a.s1 = 0; -#endif -#if defined (REPEAT_S3) -#define repeat_s(a) a.s0 = 0;a.s1 = 0;a.s2 = 0; -#endif -#if defined (REPEAT_S4) -#define repeat_s(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0; -#endif -#if defined (REPEAT_S5) -#define repeat_s(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0; -#endif -#if defined (REPEAT_S6) -#define repeat_s(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0;a.s5 = 0; -#endif -#if defined (REPEAT_S7) -#define repeat_s(a) a.s0 = 0;a.s1 = 0;a.s2 = 0;a.s3 = 0;a.s4 = 0;a.s5 = 0;a.s6 = 0; -#endif +/**************************************Count NonZero**************************************/ -#if defined (REPEAT_E0) -#define repeat_e(a) a = a; -#endif -#if defined (REPEAT_E1) -#define repeat_e(a) a.s7 = 0; -#endif -#if defined (REPEAT_E2) -#define repeat_e(a) a.s7 = 0;a.s6 = 0; -#endif -#if defined (REPEAT_E3) -#define repeat_e(a) a.s7 = 0;a.s6 = 0;a.s5 = 0; -#endif -#if defined (REPEAT_E4) -#define repeat_e(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0; -#endif -#if defined (REPEAT_E5) -#define repeat_e(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0; -#endif -#if defined (REPEAT_E6) -#define repeat_e(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0; -#endif -#if defined (REPEAT_E7) -#define repeat_e(a) a.s7 = 0;a.s6 = 0;a.s5 = 0;a.s4 = 0;a.s3 = 0;a.s2 = 0;a.s1 = 0; -#endif +__kernel void arithm_op_nonzero(int cols, int invalid_cols, int offset, int elemnum, int groupnum, + __global srcT *src, __global dstT *dst) +{ + unsigned int lid = get_local_id(0); + unsigned int gid = get_group_id(0); + unsigned int id = get_global_id(0); -#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics:enable -#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable + unsigned int idx = offset + id + (id / cols) * invalid_cols; + __local dstT localmem_nonzero[128]; + dstT nonzero = (dstT)(0); + srcT zero = (srcT)(0), one = (srcT)(1); -/**************************************Count NonZero**************************************/ -__kernel void arithm_op_nonzero (int cols,int invalid_cols,int offset,int elemnum,int groupnum, - __global VEC_TYPE *src, __global int8 *dst) -{ - unsigned int lid = get_local_id(0); - unsigned int gid = get_group_id(0); - unsigned int id = get_global_id(0); - unsigned int idx = offset + id + (id / cols) * invalid_cols; - __local int8 localmem_nonzero[128]; - int8 nonzero; - VEC_TYPE zero=0,one=1,temp; - if(id < elemnum) - { - temp = src[idx]; - if(id % cols == 0 ) - { - repeat_s(temp); - } - if(id % cols == cols - 1) - { - repeat_e(temp); - } - nonzero = convert_int8(temp == zero ? zero:one); - } - else - { - nonzero = 0; - } - for(id=id + (groupnum << 8); id < elemnum;id = id + (groupnum << 8)) - { - idx = offset + id + (id / cols) * invalid_cols; - temp = src[idx]; - if(id % cols == 0 ) - { - repeat_s(temp); - } - if(id % cols == cols - 1) - { - repeat_e(temp); - } - nonzero = nonzero + convert_int8(temp == zero ? zero:one); - } - if(lid > 127) - { - localmem_nonzero[lid - 128] = nonzero; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 128) - { - localmem_nonzero[lid] = nonzero + localmem_nonzero[lid]; - } - barrier(CLK_LOCAL_MEM_FENCE); - for(int lsize = 64; lsize > 0; lsize >>= 1) - { - if(lid < lsize) - { + for (int grain = groupnum << 8; id < elemnum; id += grain) + { + idx = offset + id + (id / cols) * invalid_cols; + nonzero += src[idx] == zero ? zero : one; + } + + if (lid > 127) + localmem_nonzero[lid - 128] = nonzero; + barrier(CLK_LOCAL_MEM_FENCE); + + if (lid < 128) + localmem_nonzero[lid] = nonzero + localmem_nonzero[lid]; + barrier(CLK_LOCAL_MEM_FENCE); + + for (int lsize = 64; lsize > 0; lsize >>= 1) + { + if (lid < lsize) + { int lid2 = lsize + lid; localmem_nonzero[lid] = localmem_nonzero[lid] + localmem_nonzero[lid2]; - } - barrier(CLK_LOCAL_MEM_FENCE); - } - if( lid == 0) - { - dst[gid] = localmem_nonzero[0]; - } + } + barrier(CLK_LOCAL_MEM_FENCE); + } + + if (lid == 0) + dst[gid] = localmem_nonzero[0]; }