From 504bc7634aeb001fe1c8bd712e2a9c2071f6ea3e Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 16 Jun 2014 13:07:39 +0400 Subject: [PATCH 1/4] Remove pre_invalid parameter --- modules/imgproc/src/opencl/integral_sum.cl | 16 ++++++++-------- modules/imgproc/src/sumpixels.cpp | 6 +++--- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index dda0e6018..728e885f4 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -63,7 +63,7 @@ #if sdepth == 4 kernel void integral_sum_cols(__global uchar4 *src, __global int *sum, - int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step) + int src_offset, int rows, int cols, int src_step, int dst_step) { int lid = get_local_id(0); int gid = get_group_id(0); @@ -122,19 +122,19 @@ kernel void integral_sum_cols(__global uchar4 *src, __global int *sum, barrier(CLK_LOCAL_MEM_FENCE); if(lid > 0 && (i+lid) <= rows) { - int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ; + int loc_s0 = gid * dst_step + i + lid - 1, loc_s1 = loc_s0 + dst_step ; lm_sum[0][bf_loc] += sum_t[0]; lm_sum[1][bf_loc] += sum_t[1]; sum_p = (__local int*)(&(lm_sum[0][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue; + if(gid * 4 + k >= cols) continue; sum[loc_s0 + k * dst_step / 4] = sum_p[k]; } sum_p = (__local int*)(&(lm_sum[1][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 4 + k + 4 >= cols + pre_invalid) break; + if(gid * 4 + k + 4 >= cols) break; sum[loc_s1 + k * dst_step / 4] = sum_p[k]; } } @@ -238,7 +238,7 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum, #elif sdepth == 5 kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, - int src_offset, int pre_invalid, int rows, int cols, int src_step, int dst_step) + int src_offset, int rows, int cols, int src_step, int dst_step) { int lid = get_local_id(0); int gid = get_group_id(0); @@ -297,19 +297,19 @@ kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, barrier(CLK_LOCAL_MEM_FENCE); if(lid > 0 && (i+lid) <= rows) { - int loc_s0 = gid * dst_step + i + lid - 1 - pre_invalid * dst_step / 4, loc_s1 = loc_s0 + dst_step ; + int loc_s0 = gid * dst_step + i + lid - 1, loc_s1 = loc_s0 + dst_step ; lm_sum[0][bf_loc] += sum_t[0]; lm_sum[1][bf_loc] += sum_t[1]; sum_p = (__local float*)(&(lm_sum[0][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 4 + k >= cols + pre_invalid || gid * 4 + k < pre_invalid) continue; + if(gid * 4 + k >= cols) continue; sum[loc_s0 + k * dst_step / 4] = sum_p[k]; } sum_p = (__local float*)(&(lm_sum[1][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 4 + k + 4 >= cols + pre_invalid) break; + if(gid * 4 + k + 4 >= cols) break; sum[loc_s1 + k * dst_step / 4] = sum_p[k]; } } diff --git a/modules/imgproc/src/sumpixels.cpp b/modules/imgproc/src/sumpixels.cpp index 111e6b67e..6b0183afa 100755 --- a/modules/imgproc/src/sumpixels.cpp +++ b/modules/imgproc/src/sumpixels.cpp @@ -254,12 +254,12 @@ static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth ) UMat src = _src.getUMat(), t_sum(t_size, sdepth), sum = _sum.getUMat(); t_sum = t_sum(Range::all(), Range(0, size.height)); - int offset = (int)src.offset / vlen, pre_invalid = (int)src.offset % vlen; - int vcols = (pre_invalid + src.cols + vlen - 1) / vlen; + int offset = (int)src.offset / vlen; + int vcols = (src.cols + vlen - 1) / vlen; int sum_offset = (int)sum.offset / vlen; k1.args(ocl::KernelArg::PtrReadOnly(src), ocl::KernelArg::PtrWriteOnly(t_sum), - offset, pre_invalid, src.rows, src.cols, (int)src.step, (int)t_sum.step); + offset, src.rows, src.cols, (int)src.step, (int)t_sum.step); size_t gt = ((vcols + 1) / 2) * 256, lt = 256; if (!k1.run(1, >, <, false)) return false; From 6550c4f682b27b8d5d6cd8543f7097aad1ec61a2 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 16 Jun 2014 15:08:15 +0400 Subject: [PATCH 2/4] Join kernel code for int and float destination types --- modules/imgproc/src/opencl/integral_sum.cl | 257 ++++----------------- 1 file changed, 46 insertions(+), 211 deletions(-) diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index 728e885f4..c9a55ac72 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -61,199 +61,34 @@ #define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS) #if sdepth == 4 - -kernel void integral_sum_cols(__global uchar4 *src, __global int *sum, - int src_offset, int rows, int cols, int src_step, int dst_step) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - int4 src_t[2], sum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local int* sum_p; - src_step = src_step >> 2; - gid = gid << 1; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid]) : 0); - src_t[1] = (i + lid < rows ? convert_int4(src[src_offset + (lid+i) * src_step + gid + 1]) : 0); - - sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - - lm_sum[1][bf_loc] = src_t[1]; - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid > 0 && (i+lid) <= rows) - { - int loc_s0 = gid * dst_step + i + lid - 1, loc_s1 = loc_s0 + dst_step ; - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local int*)(&(lm_sum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k >= cols) continue; - sum[loc_s0 + k * dst_step / 4] = sum_p[k]; - } - sum_p = (__local int*)(&(lm_sum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 4 + k + 4 >= cols) break; - sum[loc_s1 + k * dst_step / 4] = sum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - -kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum, - int rows, int cols, int src_step, int sum_step, int sum_offset) -{ - int lid = get_local_id(0); - int gid = get_group_id(0); - int4 src_t[2], sum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local int *sum_p; - src_step = src_step >> 4; - for(int i = 0; i < rows; i =i + LSIZE_1) - { - src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : 0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : 0; - - sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); - barrier(CLK_LOCAL_MEM_FENCE); - - int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; - - lm_sum[1][bf_loc] = src_t[1]; - - int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) - { - barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - } - offset <<= 1; - } - barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) - { - lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; - } - for(int d = 1; d < LSIZE; d <<= 1) - { - barrier(CLK_LOCAL_MEM_FENCE); - offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; - ai += GET_CONFLICT_OFFSET(ai); - bi += GET_CONFLICT_OFFSET(bi); - - if((lid & 127) < d) - { - lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; - lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - if(gid == 0 && (i + lid) <= rows) - { - sum[sum_offset + i + lid] = 0; - } - if(i + lid == 0) - { - int loc0 = gid * 2 * sum_step; - for(int k = 1; k <= 8; k++) - { - if(gid * 8 + k > cols) break; - sum[sum_offset + loc0 + k * sum_step / 4] = 0; - } - } - - if(lid > 0 && (i+lid) <= rows) - { - int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; - lm_sum[0][bf_loc] += sum_t[0]; - lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local int*)(&(lm_sum[0][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 8 + k >= cols) break; - sum[loc_s0 + k * sum_step / 4] = sum_p[k]; - } - sum_p = (__local int*)(&(lm_sum[1][bf_loc])); - for(int k = 0; k < 4; k++) - { - if(gid * 8 + 4 + k >= cols) break; - sum[loc_s1 + k * sum_step / 4] = sum_p[k]; - } - } - barrier(CLK_LOCAL_MEM_FENCE); - } -} - +#define sumT int +#define vecSumT int4 +#define convertToSum4 convert_int4 #elif sdepth == 5 +#define sumT float +#define vecSumT float4 +#define convertToSum4 convert_float4 +#endif -kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, + +kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, int src_offset, int rows, int cols, int src_step, int dst_step) { + sumT *sum = (sumT *)sum_ptr; int lid = get_local_id(0); int gid = get_group_id(0); - float4 src_t[2], sum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local float* sum_p; + vecSumT src_t[2], sum_t[2]; + __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE]; + __local sumT* sum_p; src_step = src_step >> 2; gid = gid << 1; - for(int i = 0; i < rows; i =i + LSIZE_1) + for (int i = 0; i < rows; i =i + LSIZE_1) { - src_t[0] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + gid]) : (float4)0); - src_t[1] = (i + lid < rows ? convert_float4(src[src_offset + (lid+i) * src_step + gid + 1]) : (float4)0); + src_t[0] = (i + lid < rows ? convertToSum4(src[mad24((lid+i), src_step, src_offset + gid)]) : (vecSumT)0); + src_t[1] = (i + lid < rows ? convertToSum4(src[mad24((lid+i), src_step, src_offset + gid + 1)]) : (vecSumT)0); - sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); + sum_t[0] = (i == 0 ? (vecSumT)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); + sum_t[1] = (i == 0 ? (vecSumT)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); barrier(CLK_LOCAL_MEM_FENCE); int bf_loc = lid + GET_CONFLICT_OFFSET(lid); @@ -262,7 +97,7 @@ kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, lm_sum[1][bf_loc] = src_t[1]; int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) + for (int d = LSIZE >> 1 ; d > 0; d>>=1) { barrier(CLK_LOCAL_MEM_FENCE); int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; @@ -276,11 +111,11 @@ kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, offset <<= 1; } barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) + if (lid < 2) { lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; } - for(int d = 1; d < LSIZE; d <<= 1) + for (int d = 1; d < LSIZE; d <<= 1) { barrier(CLK_LOCAL_MEM_FENCE); offset >>= 1; @@ -295,19 +130,19 @@ kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, } } barrier(CLK_LOCAL_MEM_FENCE); - if(lid > 0 && (i+lid) <= rows) + if (lid > 0 && (i+lid) <= rows) { - int loc_s0 = gid * dst_step + i + lid - 1, loc_s1 = loc_s0 + dst_step ; + int loc_s0 = mad24(gid, dst_step, i + lid - 1), loc_s1 = loc_s0 + dst_step; lm_sum[0][bf_loc] += sum_t[0]; lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local float*)(&(lm_sum[0][bf_loc])); - for(int k = 0; k < 4; k++) + sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); + for (int k = 0; k < 4; k++) { if(gid * 4 + k >= cols) continue; sum[loc_s0 + k * dst_step / 4] = sum_p[k]; } - sum_p = (__local float*)(&(lm_sum[1][bf_loc])); - for(int k = 0; k < 4; k++) + sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); + for (int k = 0; k < 4; k++) { if(gid * 4 + k + 4 >= cols) break; sum[loc_s1 + k * dst_step / 4] = sum_p[k]; @@ -317,22 +152,24 @@ kernel void integral_sum_cols(__global uchar4 *src, __global float *sum, } } -kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, +kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_ptr, int rows, int cols, int src_step, int sum_step, int sum_offset) { + vecSumT *srcsum = (vecSumT *)srcsum_ptr; + sumT *sum = (sumT *)sum_ptr; int lid = get_local_id(0); int gid = get_group_id(0); - float4 src_t[2], sum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local float *sum_p; + vecSumT src_t[2], sum_t[2]; + __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE]; + __local sumT *sum_p; src_step = src_step >> 4; - for(int i = 0; i < rows; i =i + LSIZE_1) + for (int i = 0; i < rows; i =i + LSIZE_1) { - src_t[0] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2] : (float4)0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : (float4)0; + src_t[0] = i + lid < rows ? srcsum[mad24((lid+i), src_step, gid * 2)] : 0; + src_t[1] = i + lid < rows ? srcsum[mad24((lid+i), src_step, gid * 2 + 1)] : 0; - sum_t[0] = (i == 0 ? (float4)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? (float4)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); + sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); + sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); barrier(CLK_LOCAL_MEM_FENCE); int bf_loc = lid + GET_CONFLICT_OFFSET(lid); @@ -341,7 +178,7 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, lm_sum[1][bf_loc] = src_t[1]; int offset = 1; - for(int d = LSIZE >> 1 ; d > 0; d>>=1) + for (int d = LSIZE >> 1 ; d > 0; d>>=1) { barrier(CLK_LOCAL_MEM_FENCE); int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; @@ -355,11 +192,11 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, offset <<= 1; } barrier(CLK_LOCAL_MEM_FENCE); - if(lid < 2) + if (lid < 2) { lm_sum[lid][LSIZE_2 + LOG_LSIZE] = 0; } - for(int d = 1; d < LSIZE; d <<= 1) + for (int d = 1; d < LSIZE; d <<= 1) { barrier(CLK_LOCAL_MEM_FENCE); offset >>= 1; @@ -374,11 +211,11 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, } } barrier(CLK_LOCAL_MEM_FENCE); - if(gid == 0 && (i + lid) <= rows) + if (gid == 0 && (i + lid) <= rows) { sum[sum_offset + i + lid] = 0; } - if(i + lid == 0) + if (i + lid == 0) { int loc0 = gid * 2 * sum_step; for(int k = 1; k <= 8; k++) @@ -388,18 +225,18 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, } } - if(lid > 0 && (i+lid) <= rows) + if (lid > 0 && (i+lid) <= rows) { int loc_s0 = sum_offset + gid * 2 * sum_step + sum_step / 4 + i + lid, loc_s1 = loc_s0 + sum_step ; lm_sum[0][bf_loc] += sum_t[0]; lm_sum[1][bf_loc] += sum_t[1]; - sum_p = (__local float*)(&(lm_sum[0][bf_loc])); + sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); for(int k = 0; k < 4; k++) { if(gid * 8 + k >= cols) break; sum[loc_s0 + k * sum_step / 4] = sum_p[k]; } - sum_p = (__local float*)(&(lm_sum[1][bf_loc])); + sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); for(int k = 0; k < 4; k++) { if(gid * 8 + 4 + k >= cols) break; @@ -409,5 +246,3 @@ kernel void integral_sum_rows(__global float4 *srcsum, __global float *sum, barrier(CLK_LOCAL_MEM_FENCE); } } - -#endif From 9bf296eeb00e0e235e5122b784711eb805614c98 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 16 Jun 2014 17:17:16 +0400 Subject: [PATCH 3/4] Small refactoring --- modules/imgproc/src/opencl/integral_sum.cl | 87 ++++++++++++++++------ 1 file changed, 64 insertions(+), 23 deletions(-) diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index c9a55ac72..a5a2ffd2a 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -82,25 +82,43 @@ kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, __local sumT* sum_p; src_step = src_step >> 2; gid = gid << 1; - for (int i = 0; i < rows; i =i + LSIZE_1) + int lid_prim = ((lid & 127) << 1) + 1; + for (int i = 0; i < rows; i += LSIZE_1) { - src_t[0] = (i + lid < rows ? convertToSum4(src[mad24((lid+i), src_step, src_offset + gid)]) : (vecSumT)0); - src_t[1] = (i + lid < rows ? convertToSum4(src[mad24((lid+i), src_step, src_offset + gid + 1)]) : (vecSumT)0); + if (i + lid < rows) + { + int src_index = mad24((lid+i), src_step, gid + src_offset); + src_t[0] = convertToSum4(src[src_index]); + src_t[1] = convertToSum4(src[src_index + 1]); + } + else + { + src_t[0] = (vecSumT)0; + src_t[1] = (vecSumT)0; + } - sum_t[0] = (i == 0 ? (vecSumT)0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? (vecSumT)0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); + if (i == 0) + { + sum_t[0] = (vecSumT)0; + sum_t[1] = (vecSumT)0; + } + else + { + sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE]; + sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE]; + } barrier(CLK_LOCAL_MEM_FENCE); int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; + lm_sum[0][bf_loc] = src_t[0]; lm_sum[1][bf_loc] = src_t[1]; int offset = 1; for (int d = LSIZE >> 1 ; d > 0; d>>=1) { barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; + int ai = offset * lid_prim - 1,bi = ai + offset; ai += GET_CONFLICT_OFFSET(ai); bi += GET_CONFLICT_OFFSET(bi); @@ -119,7 +137,7 @@ kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, { barrier(CLK_LOCAL_MEM_FENCE); offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; + int ai = offset * lid_prim - 1,bi = ai + offset; ai += GET_CONFLICT_OFFSET(ai); bi += GET_CONFLICT_OFFSET(bi); @@ -138,13 +156,15 @@ kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); for (int k = 0; k < 4; k++) { - if(gid * 4 + k >= cols) continue; + if (gid * 4 + k >= cols) + break; sum[loc_s0 + k * dst_step / 4] = sum_p[k]; } sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); for (int k = 0; k < 4; k++) { - if(gid * 4 + k + 4 >= cols) break; + if (gid * 4 + k + 4 >= cols) + break; sum[loc_s1 + k * dst_step / 4] = sum_p[k]; } } @@ -152,6 +172,7 @@ kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, } } + kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_ptr, int rows, int cols, int src_step, int sum_step, int sum_offset) { @@ -163,25 +184,42 @@ kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_pt __local vecSumT lm_sum[2][LSIZE + LOG_LSIZE]; __local sumT *sum_p; src_step = src_step >> 4; - for (int i = 0; i < rows; i =i + LSIZE_1) + int lid_prim = ((lid & 127) << 1) + 1; + for (int i = 0; i < rows; i += LSIZE_1) { - src_t[0] = i + lid < rows ? srcsum[mad24((lid+i), src_step, gid * 2)] : 0; - src_t[1] = i + lid < rows ? srcsum[mad24((lid+i), src_step, gid * 2 + 1)] : 0; - - sum_t[0] = (i == 0 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); + if (i + lid < rows) + { + int sum_idx = mad24(lid + i, src_step, gid * 2); + src_t[0] = srcsum[sum_idx]; + src_t[1] = srcsum[sum_idx + 1]; + } + else + { + src_t[0] = 0; + src_t[1] = 0; + } + if (i == 0) + { + sum_t[0] = 0; + sum_t[1] = 0; + } + else + { + sum_t[0] = lm_sum[0][LSIZE_2 + LOG_LSIZE]; + sum_t[1] = lm_sum[1][LSIZE_2 + LOG_LSIZE]; + } barrier(CLK_LOCAL_MEM_FENCE); int bf_loc = lid + GET_CONFLICT_OFFSET(lid); - lm_sum[0][bf_loc] = src_t[0]; + lm_sum[0][bf_loc] = src_t[0]; lm_sum[1][bf_loc] = src_t[1]; int offset = 1; for (int d = LSIZE >> 1 ; d > 0; d>>=1) { barrier(CLK_LOCAL_MEM_FENCE); - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; + int ai = offset * lid_prim - 1, bi = ai + offset; ai += GET_CONFLICT_OFFSET(ai); bi += GET_CONFLICT_OFFSET(bi); @@ -200,11 +238,11 @@ kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_pt { barrier(CLK_LOCAL_MEM_FENCE); offset >>= 1; - int ai = offset * (((lid & 127)<<1) +1) - 1,bi = ai + offset; + int ai = offset * lid_prim - 1,bi = ai + offset; ai += GET_CONFLICT_OFFSET(ai); bi += GET_CONFLICT_OFFSET(bi); - if((lid & 127) < d) + if ((lid & 127) < d) { lm_sum[lid >> 7][bi] += lm_sum[lid >> 7][ai]; lm_sum[lid >> 7][ai] = lm_sum[lid >> 7][bi] - lm_sum[lid >> 7][ai]; @@ -220,7 +258,8 @@ kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_pt int loc0 = gid * 2 * sum_step; for(int k = 1; k <= 8; k++) { - if(gid * 8 + k > cols) break; + if (gid * 8 + k > cols) + break; sum[sum_offset + loc0 + k * sum_step / 4] = 0; } } @@ -233,13 +272,15 @@ kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_pt sum_p = (__local sumT*)(&(lm_sum[0][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 8 + k >= cols) break; + if (gid * 8 + k >= cols) + break; sum[loc_s0 + k * sum_step / 4] = sum_p[k]; } sum_p = (__local sumT*)(&(lm_sum[1][bf_loc])); for(int k = 0; k < 4; k++) { - if(gid * 8 + 4 + k >= cols) break; + if (gid * 8 + 4 + k >= cols) + break; sum[loc_s1 + k * sum_step / 4] = sum_p[k]; } } From 606df0469aa990590d78a7b66444c054d943496e Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 16 Jun 2014 18:14:05 +0400 Subject: [PATCH 4/4] Fix pointer conversion --- modules/imgproc/src/opencl/integral_sum.cl | 10 +++++----- modules/imgproc/src/sumpixels.cpp | 2 +- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index a5a2ffd2a..333c7121c 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -71,10 +71,10 @@ #endif -kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, +kernel void integral_sum_cols(__global const uchar4 *src, __global uchar *sum_ptr, int src_offset, int rows, int cols, int src_step, int dst_step) { - sumT *sum = (sumT *)sum_ptr; + __global sumT *sum = (__global sumT *)sum_ptr; int lid = get_local_id(0); int gid = get_group_id(0); vecSumT src_t[2], sum_t[2]; @@ -173,11 +173,11 @@ kernel void integral_sum_cols(__global uchar4 *src, __global uchar *sum_ptr, } -kernel void integral_sum_rows(__global uchar *srcsum_ptr, __global uchar *sum_ptr, +kernel void integral_sum_rows(__global const uchar *srcsum_ptr, __global uchar *sum_ptr, int rows, int cols, int src_step, int sum_step, int sum_offset) { - vecSumT *srcsum = (vecSumT *)srcsum_ptr; - sumT *sum = (sumT *)sum_ptr; + __global const vecSumT *srcsum = (__global const vecSumT *)srcsum_ptr; + __global sumT *sum = (__global sumT *)sum_ptr; int lid = get_local_id(0); int gid = get_group_id(0); vecSumT src_t[2], sum_t[2]; diff --git a/modules/imgproc/src/sumpixels.cpp b/modules/imgproc/src/sumpixels.cpp index 6b0183afa..1d246ec7b 100755 --- a/modules/imgproc/src/sumpixels.cpp +++ b/modules/imgproc/src/sumpixels.cpp @@ -266,7 +266,7 @@ static bool ocl_integral( InputArray _src, OutputArray _sum, int sdepth ) ocl::Kernel k2("integral_sum_rows", ocl::imgproc::integral_sum_oclsrc, format("-D sdepth=%d", sdepth)); - k2.args(ocl::KernelArg::PtrReadWrite(t_sum), ocl::KernelArg::PtrWriteOnly(sum), + k2.args(ocl::KernelArg::PtrReadOnly(t_sum), ocl::KernelArg::PtrWriteOnly(sum), t_sum.rows, t_sum.cols, (int)t_sum.step, (int)sum.step, sum_offset); size_t gt2 = t_sum.cols * 32, lt2 = 256;