[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