[FFmpeg-devel] [PATCH] libavfilter/unsharp: add opencl unsharp filter

Stefano Sabatini stefasab at gmail.com
Thu Apr 25 12:52:41 CEST 2013


On date Thursday 2013-04-25 11:00:46 +0800, Wei Gao encoded:
> Hi,
> Modified according the comments
> 
> Thanks
> 
> 
> 2013/4/25 Stefano Sabatini <stefasab at gmail.com>
> 
> > On date Wednesday 2013-04-24 12:48:23 +0800, Wei Gao encoded:
> > > Hi
> > >
> >
> > [...]
> >
> > LGTM otherwise, thanks.
> > --
> > FFmpeg = Free & Freak Mysterious Peaceful Experimenting Gangster
> > _______________________________________________
> > ffmpeg-devel mailing list
> > ffmpeg-devel at ffmpeg.org
> > http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
> >

> From 758d5f4146181cbf5849937bb20c65aa3c9d9a41 Mon Sep 17 00:00:00 2001
> From: highgod0401 <highgod0401 at gmail.com>
> Date: Thu, 25 Apr 2013 10:47:28 +0800
> Subject: [PATCH 1/2] lavu/opencl:add opencl set param function
> 
> ---
>  libavutil/Makefile          |  2 +-
>  libavutil/opencl_internal.c | 59 +++++++++++++++++++++++++++++++++++++++++++++
>  libavutil/opencl_internal.h | 33 +++++++++++++++++++++++++
>  3 files changed, 93 insertions(+), 1 deletion(-)
>  create mode 100644 libavutil/opencl_internal.c
>  create mode 100644 libavutil/opencl_internal.h
> 
> diff --git a/libavutil/Makefile b/libavutil/Makefile
> index 33f82ed..e14d5a4 100644
> --- a/libavutil/Makefile
> +++ b/libavutil/Makefile
> @@ -108,7 +108,7 @@ OBJS = adler32.o                                                        \
>         xtea.o                                                           \
>  
>  OBJS-$(CONFIG_LZO)                      += lzo.o
> -OBJS-$(CONFIG_OPENCL)                   += opencl.o
> +OBJS-$(CONFIG_OPENCL)                   += opencl.o opencl_internal.o
>  
>  OBJS += $(COMPAT_OBJS:%=../compat/%)
>  
> diff --git a/libavutil/opencl_internal.c b/libavutil/opencl_internal.c
> new file mode 100644
> index 0000000..f15c934
> --- /dev/null
> +++ b/libavutil/opencl_internal.c
> @@ -0,0 +1,59 @@
> +/*
> + * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
> + * Copyright (C) 2012 Li   Cao <li at multicorewareinc.com>
> + * Copyright (C) 2012 Wei  Gao <weigao at multicorewareinc.com>
> + *
> + * This file is part of FFmpeg.
> + *
> + * FFmpeg is free software; you can redistribute it and/or
> + * modify it under the terms of the GNU Lesser General Public
> + * License as published by the Free Software Foundation; either
> + * version 2.1 of the License, or (at your option) any later version.
> + *
> + * FFmpeg is distributed in the hope that it will be useful,
> + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
> + * Lesser General Public License for more details.
> + *
> + * You should have received a copy of the GNU Lesser General Public
> + * License along with FFmpeg; if not, write to the Free Software
> + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
> + */
> +
> +#include "opencl_internal.h"
> +#include "libavutil/log.h"
> +
> +int ff_opencl_set_parameter(FFOpenclParam *opencl_param, ...)
> +{
> +    int ret = 0;
> +    va_list arg_ptr;
> +    void *param;
> +    size_t param_size;
> +    cl_int status;
> +    if (!opencl_param->kernel) {
> +        av_log(opencl_param->ctx, AV_LOG_ERROR, "OpenCL kernel must be set\n");
> +        return AVERROR(EINVAL);
> +    }
> +    va_start(arg_ptr, opencl_param);
> +    do {
> +        param = va_arg(arg_ptr, void *);
> +        if (!param)
> +            break;
> +        param_size = va_arg(arg_ptr, size_t);
> +        if (!param_size) {
> +            av_log(opencl_param->ctx, AV_LOG_ERROR, "Parameter size must not be 0\n");
> +            ret = AVERROR(EINVAL);
> +            goto end;
> +        }
> +        status = clSetKernelArg(opencl_param->kernel, opencl_param->param_num, param_size, param);
> +        if (status != CL_SUCCESS) {
> +            av_log(opencl_param->ctx, AV_LOG_ERROR, "Cannot set kernel argument: %d\n", status);
> +            ret = AVERROR_EXTERNAL;
> +            goto end;
> +        }
> +        opencl_param->param_num++;
> +    } while (param && param_size);
> +end:
> +    va_end(arg_ptr);
> +    return ret;
> +}
> diff --git a/libavutil/opencl_internal.h b/libavutil/opencl_internal.h
> new file mode 100644
> index 0000000..34b39a0
> --- /dev/null
> +++ b/libavutil/opencl_internal.h
> @@ -0,0 +1,33 @@
> +/*
> + * Copyright (C) 2012 Peng Gao <peng at multicorewareinc.com>
> + * Copyright (C) 2012 Li   Cao <li at multicorewareinc.com>
> + * Copyright (C) 2012 Wei  Gao <weigao at multicorewareinc.com>
> + *
> + * This file is part of FFmpeg.
> + *
> + * FFmpeg is free software; you can redistribute it and/or
> + * modify it under the terms of the GNU Lesser General Public
> + * License as published by the Free Software Foundation; either
> + * version 2.1 of the License, or (at your option) any later version.
> + *
> + * FFmpeg is distributed in the hope that it will be useful,
> + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
> + * Lesser General Public License for more details.
> + *
> + * You should have received a copy of the GNU Lesser General Public
> + * License along with FFmpeg; if not, write to the Free Software
> + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
> + */
> +
> +#include "opencl.h"
> +
> +#define FF_OPENCL_PARAM_INFO(a) ((void*)(&(a))), (sizeof(a))
> +
> +typedef struct {
> +    cl_kernel kernel;
> +    int param_num;
> +    void *ctx;
> +} FFOpenclParam;
> +
> +int ff_opencl_set_parameter(FFOpenclParam *opencl_param, ...);
> -- 
> 1.7.11.msysgit.1

Still LGTM.

> From 134ea17f55a872c085737d3000d0ef23d267ad56 Mon Sep 17 00:00:00 2001
> From: highgod0401 <highgod0401 at gmail.com>
> Date: Thu, 25 Apr 2013 10:58:26 +0800
> Subject: [PATCH 2/2] lavfi/unsharp: add opencl unsharp filter
> 
> ---
>  doc/filters.texi                |   5 +
>  libavfilter/Makefile            |   2 +-
>  libavfilter/opencl_allkernels.c |   2 +
>  libavfilter/unsharp.h           |  78 +++++++++++
>  libavfilter/unsharp_kernel.h    | 132 +++++++++++++++++++
>  libavfilter/unsharp_opencl.c    | 282 ++++++++++++++++++++++++++++++++++++++++
>  libavfilter/unsharp_opencl.h    |  34 +++++
>  libavfilter/version.h           |   2 +-
>  libavfilter/vf_unsharp.c        |  88 ++++++++-----
>  9 files changed, 587 insertions(+), 38 deletions(-)
>  create mode 100644 libavfilter/unsharp.h
>  create mode 100644 libavfilter/unsharp_kernel.h
>  create mode 100644 libavfilter/unsharp_opencl.c
>  create mode 100644 libavfilter/unsharp_opencl.h
[...]
> diff --git a/libavfilter/unsharp_kernel.h b/libavfilter/unsharp_kernel.h
> new file mode 100644
> index 0000000..c1de11d
> --- /dev/null
> +++ b/libavfilter/unsharp_kernel.h
> @@ -0,0 +1,132 @@
> +/*
> + * Copyright (C) 2013 Wei Gao <weigao at multicorewareinc.com>
> + *
> + * This file is part of FFmpeg.
> + *
> + * FFmpeg is free software; you can redistribute it and/or
> + * modify it under the terms of the GNU Lesser General Public
> + * License as published by the Free Software Foundation; either
> + * version 2.1 of the License, or (at your option) any later version.
> + *
> + * FFmpeg is distributed in the hope that it will be useful,
> + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
> + * Lesser General Public License for more details.
> + *
> + * You should have received a copy of the GNU Lesser General Public
> + * License along with FFmpeg; if not, write to the Free Software
> + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
> + */
> +
> +#ifndef AVFILTER_UNSHARP_KERNEL_H
> +#define AVFILTER_UNSHARP_KERNEL_H
> +
> +#include "libavutil/opencl.h"
> +
> +const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
> +kernel void unsharp(global  unsigned char *src,
> +                    global  unsigned char *dst,
> +                    const global  unsigned int *mask_lu,
> +                    const global  unsigned int *mask_ch,
> +                    int amount_lu,
> +                    int amount_ch,
> +                    int step_x_lu,
> +                    int step_y_lu,
> +                    int step_x_ch,
> +                    int step_y_ch,
> +                    int scalebits_lu,
> +                    int scalebits_ch,
> +                    int halfscale_lu,
> +                    int halfscale_ch,
> +                    int src_stride_lu,
> +                    int src_stride_ch,
> +                    int dst_stride_lu,
> +                    int dst_stride_ch,
> +                    int height,
> +                    int width,
> +                    int ch,
> +                    int cw)
> +{
> +    global unsigned char *dst_y = dst;
> +    global unsigned char *dst_u = dst_y + height * dst_stride_lu;
> +    global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
> +
> +    global unsigned char *src_y = src;
> +    global unsigned char *src_u = src_y + height * src_stride_lu;
> +    global unsigned char *src_v = src_u + ch * src_stride_ch;
> +
> +    global unsigned char *temp_dst;
> +    global unsigned char *temp_src;
> +    const global unsigned int *temp_mask;
> +    int global_id = get_global_id(0);
> +    int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
> +        temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
> +    if (global_id < width * height) {
> +        y = global_id / width;
> +        x = global_id % width;
> +        temp_dst = dst_y;
> +        temp_src = src_y;
> +        temp_src_stride = src_stride_lu;
> +        temp_dst_stride = dst_stride_lu;
> +        temp_height = height;
> +        temp_width = width;
> +        temp_steps_x = step_x_lu;
> +        temp_steps_y = step_y_lu;
> +        temp_mask = mask_lu;
> +        temp_amount = amount_lu;
> +        temp_scalebits = scalebits_lu;
> +        temp_halfscale = halfscale_lu;
> +    } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
> +        y = (global_id - width * height) / cw;
> +        x = (global_id - width * height) % cw;
> +        temp_dst = dst_u;
> +        temp_src = src_u;
> +        temp_src_stride = src_stride_ch;
> +        temp_dst_stride = dst_stride_ch;
> +        temp_height = ch;
> +        temp_width = cw;
> +        temp_steps_x = step_x_ch;
> +        temp_steps_y = step_y_ch;
> +        temp_mask = mask_ch;
> +        temp_amount = amount_ch;
> +        temp_scalebits = scalebits_ch;
> +        temp_halfscale = halfscale_ch;
> +    } else {
> +        y = (global_id - width * height - ch * cw) / cw;
> +        x = (global_id - width * height - ch * cw) % cw;
> +        temp_dst = dst_v;
> +        temp_src = src_v;
> +        temp_src_stride = src_stride_ch;
> +        temp_dst_stride = dst_stride_ch;
> +        temp_height = ch;
> +        temp_width = cw;
> +        temp_steps_x = step_x_ch;
> +        temp_steps_y = step_y_ch;
> +        temp_mask = mask_ch;
> +        temp_amount = amount_ch;
> +        temp_scalebits = scalebits_ch;
> +        temp_halfscale = halfscale_ch;
> +    }
> +    if (temp_amount) {
> +        sum = 0;
> +        for (j = 0; j <= 2 * temp_steps_y; j++) {
> +            idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
> +            for (i = 0; i <= 2 * temp_steps_x; i++) {
> +                idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
> +                sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
> +            }
> +        }
> +        temp = (int)temp_src[x + y * temp_src_stride];
> +        res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);

> +        if (res & (~0xFF))
> +            temp_dst[x + y * temp_dst_stride] = (-res) >> 31;
> +        else
> +            temp_dst[x + y * temp_dst_stride] = res;

Add a note like // clip res in the 0-255 range, or even better create
an inline function for this for readability's sake.

[...]
> diff --git a/libavfilter/version.h b/libavfilter/version.h
> index ba1ec0b..e86b07d 100644
> --- a/libavfilter/version.h
> +++ b/libavfilter/version.h
> @@ -29,7 +29,7 @@
>  #include "libavutil/avutil.h"
>  
>  #define LIBAVFILTER_VERSION_MAJOR  3
> -#define LIBAVFILTER_VERSION_MINOR  56
> +#define LIBAVFILTER_VERSION_MINOR  57
>  #define LIBAVFILTER_VERSION_MICRO 103

A *micro* bump is probably better (since there is no new public symbol
define).

[...]

No more comments from me, thanks.
-- 
FFmpeg = Frightening and Fierce Merciful Ponderous Epic Gymnast


More information about the ffmpeg-devel mailing list