This is a better way to accomplish the box-blur cl operation by using ocl's

local memory from the opencv source code. It use the local shared memory to
reduce global memory access, which significantly reduces the kernel's processing
time by 70 percent compared to the original one. Because of the barriers and
local worksize limitation, processing with a radius larger than 110 becomes
slower than original algorithm, so I keep the original kernels in order to deal
with box-blur with radius larger than 110.
All the tests are based on Intel Beginet and Intel IvyBridge CPU and GPU.

v2:add kernel attribute to restrict the local size to (256,1,1).

Signed-off-by: Yongjia Zhang<yongjia.zhang@intel.com>

2 files changed, 136 insertions, 0 deletions

diff --git a/opencl/box-blur.cl b/opencl/box-blur.cl
index e99bea47..5011fa35 100644
--- a/opencl/box-blur.cl
+++ b/opencl/box-blur.cl
@@ -43,3 +43,71 @@ __kernel void kernel_blur_ver (__global const float4     *aux,
       out[out_index] = mean / (float)(2 * radius + 1);
     }
 }
+
+__kernel
+__attibute__((reqd_work_group_size(256,1,1)))
+void kernel_box_blur_fast(const __global float4 *in,
+                                   __global float4 *out,
+                                   __local float4 *column_sum,
+                                   const int width,
+                                   const int height,
+                                   const int radius,
+                                   const int size)
+{
+       const int local_id0 = get_local_id(0);
+       const int twice_radius = 2 * radius;
+       const int in_width = twice_radius + width;
+       const int in_height = twice_radius + height;
+       const float4 area = (float4)( (twice_radius+1) * (twice_radius+1) );
+       int column_index_start,column_index_end;
+       int y = get_global_id(1) * size;
+       const int out_x = get_group_id(0)
+               * ( get_local_size(0) - twice_radius ) + local_id0 - radius;
+       const int in_x = out_x + radius;
+       int tmp_size = size;
+       int tmp_index = 0;
+       float4 tmp_sum = (float4)0.0f;
+       float4 total_sum = (float4)0.0f;
+       if( in_x < in_width )
+       {
+               column_index_start = y;
+               column_index_end = y + twice_radius;
+               for( int i=0; i<twice_radius+1; ++i )
+                       tmp_sum+=in[(y+i)*in_width+in_x];
+               column_sum[local_id0] = tmp_sum;
+       }
+
+       barrier( CLK_LOCAL_MEM_FENCE );
+
+       while(1)
+       {
+               if( out_x < width )
+               {
+                       if( local_id0 >= radius
+                                 && local_id0 < get_local_size(0) - radius )
+                       {
+                               total_sum = (float4)0.0f;
+                               for( int i=0; i<twice_radius+1; ++i )
+                                  total_sum += column_sum[local_id0-radius+i];
+                               out[y*width+out_x] = total_sum/area;
+                       }
+               }
+               if( --tmp_size ==0 || y == height - 1 )
+                       break;
+
+               barrier( CLK_LOCAL_MEM_FENCE );
+
+               ++y;
+               if( in_x < in_width )
+               {
+                       tmp_sum = column_sum[local_id0];
+                       tmp_sum -= in[(column_index_start)*in_width+in_x];
+                       tmp_sum += in[(column_index_end+1)*in_width+in_x];
+                       ++column_index_start;
+                       ++column_index_end;
+                       column_sum[local_id0] = tmp_sum;
+               }
+
+               barrier( CLK_LOCAL_MEM_FENCE );
+       }
+}
diff --git a/opencl/box-blur.cl.h b/opencl/box-blur.cl.h
index bfed601d..e4585eca 100644
--- a/opencl/box-blur.cl.h
+++ b/opencl/box-blur.cl.h
@@ -44,4 +44,72 @@ static const char* box_blur_cl_source =
 "      out[out_index] = mean / (float)(2 * radius + 1);                        \n"
 "    }                                                                         \n"
 "}                                                                             \n"
+"                                                                              \n"
+"__kernel                                                                      \n"
+"__attribute__((reqd_work_group_size(256,1,1)))                                \n"
+"void kernel_box_blur_fast(const __global float4 *in,                          \n"
+"                                   __global float4 *out,                      \n"
+"                                   __local float4 *column_sum,                \n"
+"                                   const int width,                           \n"
+"                                   const int height,                          \n"
+"                                   const int radius,                          \n"
+"                                   const int size)                            \n"
+"{                                                                             \n"
+"    const int local_id0 = get_local_id(0);                                    \n"
+"    const int twice_radius = 2 * radius;                                      \n"
+"    const int in_width = twice_radius + width;                                \n"
+"    const int in_height = twice_radius + height;                              \n"
+"    const float4 area = (float4)( (twice_radius+1) * (twice_radius+1) );      \n"
+"    int column_index_start,column_index_end;                                  \n"
+"    int y = get_global_id(1) * size;                                          \n"
+"    const int out_x = get_group_id(0)                                         \n"
+"               * ( get_local_size(0) - twice_radius ) + local_id0 - radius;   \n"
+"    const int in_x = out_x + radius;                                          \n"
+"    int tmp_size = size;                                                      \n"
+"    int tmp_index = 0;                                                        \n"
+"    float4 tmp_sum = (float4)0.0f;                                            \n"
+"    float4 total_sum = (float4)0.0f;                                          \n"
+"    if( in_x < in_width )                                                     \n"
+"    {                                                                         \n"
+"        column_index_start = y;                                               \n"
+"        column_index_end = y + twice_radius;                                  \n"
+"        for( int i=0; i<twice_radius+1; ++i )                                 \n"
+"            tmp_sum+=in[(y+i)*in_width+in_x];                                 \n"
+"        column_sum[local_id0] = tmp_sum;                                      \n"
+"    }                                                                         \n"
+"                                                                              \n"
+"    barrier( CLK_LOCAL_MEM_FENCE );                                           \n"
+"                                                                              \n"
+"    while(1)                                                                  \n"
+"    {                                                                         \n"
+"        if( out_x < width )                                                   \n"
+"        {                                                                     \n"
+"            if( local_id0 >= radius                                           \n"
+"                                 && local_id0 < get_local_size(0) - radius )  \n"
+"            {                                                                 \n"
+"                total_sum = (float4)0.0f;                                     \n"
+"                for( int i=0; i<twice_radius+1; ++i )                         \n"
+"                   total_sum += column_sum[local_id0-radius+i];               \n"
+"                out[y*width+out_x] = total_sum/area;                          \n"
+"            }                                                                 \n"
+"        }                                                                     \n"
+"        if( --tmp_size ==0 || y == height - 1 )                               \n"
+"            break;                                                            \n"
+"                                                                              \n"
+"        barrier( CLK_LOCAL_MEM_FENCE );                                       \n"
+"                                                                              \n"
+"        ++y;                                                                  \n"
+"        if( in_x < in_width )                                                 \n"
+"        {                                                                     \n"
+"            tmp_sum = column_sum[local_id0];                                  \n"
+"            tmp_sum -= in[(column_index_start)*in_width+in_x];                \n"
+"            tmp_sum += in[(column_index_end+1)*in_width+in_x];                \n"
+"            ++column_index_start;                                             \n"
+"            ++column_index_end;                                               \n"
+"            column_sum[local_id0] = tmp_sum;                                  \n"
+"        }                                                                     \n"
+"                                                                              \n"
+"        barrier( CLK_LOCAL_MEM_FENCE );                                       \n"
+"    }                                                                         \n"
+"}                                                                             \n"
 ;