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

refactoring

Browse Source
This commit is contained in:
Ilya Lavrenov
2013-12-03 00:41:07 +04:00
parent 49e729f89d
commit 49474903cb
5 changed files with 51 additions and 39 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
modules/core/src/opencl/reduce.cl
View File

@@ -66,7 +66,7 @@
dstT temp = convertToDT(src[0]); \
FUNC(accumulator, temp)
#define REDUCE_LOCAL_1 \
localmem[lid] += accumulator
localmem[lid - WGS2_ALIGNED] += accumulator
#define REDUCE_LOCAL_2 \
localmem[lid] += localmem[lid2]
@@ -78,7 +78,7 @@
#define REDUCE_GLOBAL \
accumulator += src[0] == zero ? zero : one
#define REDUCE_LOCAL_1 \
localmem[lid] += accumulator
localmem[lid - WGS2_ALIGNED] += accumulator
#define REDUCE_LOCAL_2 \
localmem[lid] += localmem[lid2]
@@ -95,10 +95,6 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
int id = get_global_id(0);
__local dstT localmem[WGS2_ALIGNED];
if (lid < WGS2_ALIGNED)
localmem[lid] = (dstT)(0);
barrier(CLK_LOCAL_MEM_FENCE);
DEFINE_ACCUMULATOR;
for (int grain = groupnum * WGS; id < total; id += grain)
@@ -108,11 +104,11 @@ __kernel void reduce(__global const uchar * srcptr, int step, int offset, int co
REDUCE_GLOBAL;
}
if (lid >= WGS2_ALIGNED)
localmem[lid - WGS2_ALIGNED] = accumulator;
if (lid < WGS2_ALIGNED)
localmem[lid] = accumulator;
barrier(CLK_LOCAL_MEM_FENCE);
if (lid < WGS2_ALIGNED)
if (lid >= WGS2_ALIGNED)
REDUCE_LOCAL_1;
barrier(CLK_LOCAL_MEM_FENCE);