Re: [PATCH] Optimize operation box-blur opencl kernel

[Date Prev][Date Next][Thread Prev][Thread Next][Date Index][Thread Index]

 



Thanks Yongjia, I think this looks like a nice contribution. I'll push
the patch soon.

Victor Oliveira

On Wed, Jan 22, 2014 at 12:21 AM, Yongjia Zhang <zhang_yong_jia@xxxxxxx> wrote:
> From: Yongjia Zhang <Zhang_Yong_jia@xxxxxxx>
>
> 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.
>
> Signed-off-by: Yongjia Zhang <yongjia.zhang@xxxxxxxxx>
> ---
>  opencl/box-blur.cl           |  66 +++++++++++++++++++++++++
>  opencl/box-blur.cl.h         |  66 +++++++++++++++++++++++++
>  operations/common/box-blur.c | 115 ++++++++++++++++++++++++++-----------------
>  3 files changed, 201 insertions(+), 46 deletions(-)
>
> diff --git a/opencl/box-blur.cl b/opencl/box-blur.cl
> index e99bea4..a1da9de 100644
> --- a/opencl/box-blur.cl
> +++ b/opencl/box-blur.cl
> @@ -43,3 +43,69 @@ __kernel void kernel_blur_ver (__global const float4     *aux,
>        out[out_index] = mean / (float)(2 * radius + 1);
>      }
>  }
> +
> +__kernel 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 bfed601..8f6aa81 100644
> --- a/opencl/box-blur.cl.h
> +++ b/opencl/box-blur.cl.h
> @@ -44,4 +44,70 @@ static const char* box_blur_cl_source =
>  "      out[out_index] = mean / (float)(2 * radius + 1);                        \n"
>  "    }                                                                         \n"
>  "}                                                                             \n"
> +"                                                                              \n"
> +"__kernel 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"
>  ;
> diff --git a/operations/common/box-blur.c b/operations/common/box-blur.c
> index afc19ea..cb77ec0 100644
> --- a/operations/common/box-blur.c
> +++ b/operations/common/box-blur.c
> @@ -180,9 +180,7 @@ static void prepare (GeglOperation *operation)
>  #include "buffer/gegl-buffer-cl-iterator.h"
>
>  #include "opencl/box-blur.cl.h"
> -
>  static GeglClRunData *cl_data = NULL;
> -
>  static gboolean
>  cl_box_blur (cl_mem                in_tex,
>               cl_mem                aux_tex,
> @@ -192,57 +190,82 @@ cl_box_blur (cl_mem                in_tex,
>               gint                  radius)
>  {
>    cl_int cl_err = 0;
> -  size_t global_ws_hor[2], global_ws_ver[2];
> -  size_t local_ws_hor[2], local_ws_ver[2];
> -
> +  size_t global_ws_hor[2], global_ws_ver[2], global_ws[2];
> +  size_t local_ws_hor[2], local_ws_ver[2], local_ws[2];
> +  size_t step_size ;
>    if (!cl_data)
>      {
> -      const char *kernel_name[] = {"kernel_blur_hor", "kernel_blur_ver", NULL};
> +      const char *kernel_name[] = { "kernel_blur_hor", "kernel_blur_ver","kernel_box_blur_fast", NULL};
>        cl_data = gegl_cl_compile_and_build (box_blur_cl_source, kernel_name);
>      }
>
>    if (!cl_data)
>      return TRUE;
> -
> -  local_ws_hor[0] = 1;
> -  local_ws_hor[1] = 256;
> -  global_ws_hor[0] = roi->height + 2 * radius;
> -  global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1];
> -
> -  local_ws_ver[0] = 1;
> -  local_ws_ver[1] = 256;
> -  global_ws_ver[0] = roi->height;
> -  global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];
> -
> -
> -  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
> -                                    sizeof(cl_mem), (void*)&in_tex,
> -                                    sizeof(cl_mem), (void*)&aux_tex,
> -                                    sizeof(cl_int), (void*)&roi->width,
> -                                    sizeof(cl_int), (void*)&radius,
> -                                    NULL);
> -  CL_CHECK;
> -
> -  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
> -                                        cl_data->kernel[0], 2,
> -                                        NULL, global_ws_hor, local_ws_hor,
> -                                        0, NULL, NULL);
> -  CL_CHECK;
> -
> -
> -  cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
> -                                    sizeof(cl_mem), (void*)&aux_tex,
> -                                    sizeof(cl_mem), (void*)&out_tex,
> -                                    sizeof(cl_int), (void*)&roi->width,
> -                                    sizeof(cl_int), (void*)&radius,
> -                                    NULL);
> -  CL_CHECK;
> -
> -  cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
> -                                        cl_data->kernel[1], 2,
> -                                        NULL, global_ws_ver, local_ws_ver,
> -                                        0, NULL, NULL);
> -  CL_CHECK;
> +  step_size = 64;
> +  local_ws[0]=256;
> +  local_ws[1]=1;
> +
> +
> +  if( radius <=110 )
> +  {
> +    global_ws[0] = (roi->width + local_ws[0] - 2 * radius - 1) / ( local_ws[0] - 2 * radius ) * local_ws[0];
> +    global_ws[1] = (roi->height + step_size - 1) / step_size;
> +    cl_err = gegl_cl_set_kernel_args(cl_data->kernel[2],
> +                                     sizeof(cl_mem), (void *)&in_tex,
> +                                     sizeof(cl_mem), (void *)&out_tex,
> +                                     sizeof(cl_float4)*local_ws[0], (void *)NULL,
> +                                     sizeof(cl_int), (void *)&roi->width,
> +                                     sizeof(cl_int), (void *)&roi->height,
> +                                     sizeof(cl_int), (void *)&radius,
> +                                     sizeof(cl_int), (void *)&step_size, NULL);
> +    CL_CHECK;
> +    cl_err = gegl_clEnqueueNDRangeKernel(gegl_cl_get_command_queue(),
> +                                         cl_data->kernel[2], 2,
> +                                         NULL, global_ws, local_ws, 0, NULL, NULL );
> +       CL_CHECK;
> +
> +  }
> +  else
> +  {
> +    local_ws_hor[0] = 1;
> +    local_ws_hor[1] = 256;
> +    global_ws_hor[0] = roi->height + 2 * radius;
> +    global_ws_hor[1] = ((roi->width + local_ws_hor[1] -1)/local_ws_hor[1]) * local_ws_hor[1];
> +
> +    local_ws_ver[0] = 1;
> +    local_ws_ver[1] = 256;
> +    global_ws_ver[0] = roi->height;
> +    global_ws_ver[1] = ((roi->width + local_ws_ver[1] -1)/local_ws_ver[1]) * local_ws_ver[1];
> +
> +
> +    cl_err = gegl_cl_set_kernel_args (cl_data->kernel[0],
> +                                      sizeof(cl_mem), (void*)&in_tex,
> +                                      sizeof(cl_mem), (void*)&aux_tex,
> +                                      sizeof(cl_int), (void*)&roi->width,
> +                                      sizeof(cl_int), (void*)&radius,
> +                                      NULL);
> +    CL_CHECK;
> +    cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
> +                                          cl_data->kernel[0], 2,
> +                                          NULL, global_ws_hor, local_ws_hor,
> +                                          0, NULL, NULL);
> +    CL_CHECK;
> +
> +
> +    cl_err = gegl_cl_set_kernel_args (cl_data->kernel[1],
> +                                      sizeof(cl_mem), (void*)&aux_tex,
> +                                      sizeof(cl_mem), (void*)&out_tex,
> +                                      sizeof(cl_int), (void*)&roi->width,
> +                                      sizeof(cl_int), (void*)&radius,
> +                                      NULL);
> +    CL_CHECK;
> +
> +    cl_err = gegl_clEnqueueNDRangeKernel (gegl_cl_get_command_queue (),
> +                                          cl_data->kernel[1], 2,
> +                                          NULL, global_ws_ver, local_ws_ver,
> +                                          0, NULL, NULL);
> +    CL_CHECK;
> +  }
>
>    return FALSE;
>
> --
> 1.8.3.2
>
>
> _______________________________________________
> gegl-developer-list mailing list
> List address:    gegl-developer-list@xxxxxxxxx
> List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list
>
_______________________________________________
gegl-developer-list mailing list
List address:    gegl-developer-list@xxxxxxxxx
List membership: https://mail.gnome.org/mailman/listinfo/gegl-developer-list





[Index of Archives]     [Yosemite News]     [Yosemite Photos]     [gtk]     [GIMP Users]     [KDE]     [Gimp's Home]     [Gimp on Windows]     [Steve's Art]

  Powered by Linux