Join kernel code for int and float destination types

diff --git a/modules/imgproc/src/opencl/integral_sum.cl b/modules/imgproc/src/opencl/integral_sum.cl
index 728e885..c9a55ac 100644 (file)
--- 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