From 6550c4f682b27b8d5d6cd8543f7097aad1ec61a2 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 16 Jun 2014 15:08:15 +0400 Subject: [PATCH] Join kernel code for int and float destination types --- modules/imgproc/src/opencl/integral_sum.cl | 255 +++++------------------------ 1 file changed, 45 insertions(+), 210 deletions(-) diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl index 728e885..c9a55ac 100644 --- a/modules/imgproc/src/opencl/integral_sum.cl +++ b/modules/imgproc/src/opencl/integral_sum.cl @@ -61,24 +61,34 @@ #define GET_CONFLICT_OFFSET(lid) ((lid) >> LOG_NUM_BANKS) #if sdepth == 4 +#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 int *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); - int4 src_t[2], sum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local int* 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_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); + 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 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); - sum_t[1] = (i == 0 ? 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); @@ -87,7 +97,7 @@ kernel void integral_sum_cols(__global uchar4 *src, __global int *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; @@ -101,11 +111,11 @@ kernel void integral_sum_cols(__global uchar4 *src, __global int *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; @@ -120,19 +130,19 @@ kernel void integral_sum_cols(__global uchar4 *src, __global int *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 int*)(&(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 int*)(&(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]; @@ -142,19 +152,21 @@ kernel void integral_sum_cols(__global uchar4 *src, __global int *sum, } } -kernel void integral_sum_rows(__global int4 *srcsum, __global int *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); - int4 src_t[2], sum_t[2]; - __local int4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local int *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] : 0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : 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 ? 0 : lm_sum[0][LSIZE_2 + LOG_LSIZE]); sum_t[1] = (i == 0 ? 0 : lm_sum[1][LSIZE_2 + LOG_LSIZE]); @@ -166,7 +178,7 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *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; @@ -180,11 +192,11 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *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; @@ -199,11 +211,11 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *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++) @@ -213,18 +225,18 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *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 int*)(&(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 int*)(&(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; @@ -234,180 +246,3 @@ kernel void integral_sum_rows(__global int4 *srcsum, __global int *sum, barrier(CLK_LOCAL_MEM_FENCE); } } - -#elif sdepth == 5 - -kernel void integral_sum_cols(__global uchar4 *src, __global float *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); - float4 src_t[2], sum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local float* 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_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); - - 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]); - 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 float*)(&(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++) - { - 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 float4 *srcsum, __global float *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); - float4 src_t[2], sum_t[2]; - __local float4 lm_sum[2][LSIZE + LOG_LSIZE]; - __local float *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] : (float4)0; - src_t[1] = i + lid < rows ? srcsum[(lid+i) * src_step + gid * 2 + 1] : (float4)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]); - 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 float*)(&(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])); - 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); - } -} - -#endif -- 2.7.4