[FFmpeg-devel] [PATCH] libavfilter/unsharp: add opencl unsharp filter
Wei Gao
highgod0401 at gmail.com
Sun Apr 21 11:10:38 CEST 2013
Hi,
Thanks for your reviewing, Stefano, some questions as follows:
2013/4/20 Stefano Sabatini <stefasab at gmail.com>
> On date Friday 2013-04-19 19:46:56 +0800, Wei Gao encoded:
> >
>
> > From 995d82ac35f374b836629a0592debc57f0a97778 Mon Sep 17 00:00:00 2001
> > From: highgod0401 <highgod0401 at gmail.com>
> > Date: Fri, 19 Apr 2013 19:43:01 +0800
> > Subject: [PATCH] add opencl unsharp filter
> >
> > ---
> > libavfilter/Makefile | 2 +-
> > libavfilter/opencl_allkernels.c | 5 +-
> > libavfilter/unsharp.h | 79 +++++++++++
> > libavfilter/unsharp_kernel.h | 134 ++++++++++++++++++
> > libavfilter/unsharp_opencl.c | 301
> ++++++++++++++++++++++++++++++++++++++++
> > libavfilter/unsharp_opencl.h | 34 +++++
> > libavfilter/vf_unsharp.c | 87 +++++++-----
>
>
> > +#endif /* AVFILTER_UNSHARP_H */
> > diff --git a/libavfilter/unsharp_kernel.h b/libavfilter/unsharp_kernel.h
> > new file mode 100644
> > index 0000000..67fa6c0
> > --- /dev/null
> > +++ b/libavfilter/unsharp_kernel.h
>
> I think this should be renamed:
> unsharp_opencl_kernel.h
>
> (same for deshake_kernel.h)
>
Could I rename the files(deshake_kernel.h and unsharp_kernel.h) in another
patch later? Because if I rename now, I have to submit a new patch for
deshake.
>
> > + int global_id = get_global_id(0);
>
> size_t global_id_size may be a better name
>
It is kernel code which will be compiled in GPU, and there seems no data
type named size_t.
>
> > + 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) {
>
>
> > +#define UNSHARP_OPENCL_CHECK(method, ...)
> \
> > + status = method(__VA_ARGS__);
> \
> > + if (status != CL_SUCCESS) {
> \
> > + av_log(ctx, AV_LOG_ERROR, "error %s %d\n", # method, status);
> \
>
> Usual note on an av_opencl_ public error API.
>
The function is used as a intra function of opencl utils. can I add the api
in another patch later? And then modify the deshake and unsharp file?
>
> > + return AVERROR_EXTERNAL;
> \
> > + }
> > +
> > +#define UNSHARP_OPENCL_SET_KERNEL_ARG(arg_ptr)
> \
> > + status =
> clSetKernelArg((kernel),(arg_no++),(sizeof(arg_ptr)),(void*)(&(arg_ptr)));
> \
> > + if (status != CL_SUCCESS) {
> \
> > + av_log(ctx, AV_LOG_ERROR, "cannot set kernel argument: %d\n",
> status ); \
> > + return AVERROR_EXTERNAL;
> \
> > + }
>
> These macros are mostly duplicated from deshake_opencl.c, so we think
> to factorize them to a common file (libavfilter/opencl_internal.h
> could be fine).
>
Yes, I agree with you, I defined them in the previous patch but there seems
some problem because the parameter status, kernel, arg_no must be declared
in the function which use the macro. Can you give me some advice? Thanks.
>
> > +
> > +static void add_mask_counter(uint32_t *dst, uint32_t *counter1,
> uint32_t *counter2, int len)
> > +{
> > + int i;
> > + for (i = 0; i < len; i++) {
> > + dst[i] = counter1[i] + counter2[i];
> > + }
> > +}
> > +
>
> > + ret = av_opencl_buffer_create(&unsharp->opencl_ctx.cl_chroma_mask,
> > + sizeof(uint32_t) * (2 *
> unsharp->chroma.steps_x + 1) * (2 * unsharp->chroma.steps_y + 1),
> > + CL_MEM_READ_ONLY, NULL);
> > + if (ret < 0)
> > + return ret;
> > + ret = generate_mask(ctx);
> > + if (ret < 0)
> > + return ret;
>
> > + unsharp->opencl_ctx.plane_num = PLANE_NUM;
>
> potentially unsafe if we add support to a different number of planes
>
where can I get the plane number? should I have to detect data0,data1,
data2? I want to add share buffer in the later patch, maybe use data4 to
pass GPU buffer, so I don't think it is a good idea.
>
> > + if (!unsharp->opencl_ctx.kernel_env.kernel) {
>
>
> > + unsharp->opencl_ctx.in_plane_size[0] = (in->linesize[0] *
> in->height);
> > + unsharp->opencl_ctx.in_plane_size[1] = (in->linesize[1] * ch);
> > + unsharp->opencl_ctx.in_plane_size[2] = (in->linesize[2] * ch);
> > + unsharp->opencl_ctx.out_plane_size[0] = (out->linesize[0] *
> out->height);
> > + unsharp->opencl_ctx.out_plane_size[1] = (out->linesize[1] * ch);
> > + unsharp->opencl_ctx.out_plane_size[2] = (out->linesize[2] * ch);
> > + unsharp->opencl_ctx.cl_inbuf_size =
> unsharp->opencl_ctx.in_plane_size[0] +
> > +
> unsharp->opencl_ctx.in_plane_size[1] +
> > +
> unsharp->opencl_ctx.in_plane_size[2];
> > + unsharp->opencl_ctx.cl_outbuf_size =
> unsharp->opencl_ctx.out_plane_size[0] +
> > +
> unsharp->opencl_ctx.out_plane_size[1] +
> > +
> unsharp->opencl_ctx.out_plane_size[2];
>
> looks like another good candidate for factorization
>
Sorry, I don't get what you mean.
>
> >
> > static void apply_unsharp( uint8_t *dst, int dst_stride,
> > const uint8_t *src, int src_stride,
> > - int width, int height, FilterParam *fp)
> > + int width, int height, UnsharpFilterParam *fp)
> > {
> > uint32_t **sc = fp->sc;
> > uint32_t sr[MAX_MATRIX_SIZE - 1], tmp1, tmp2;
> > @@ -131,7 +107,27 @@ static void apply_unsharp( uint8_t *dst, int
> dst_stride,
> > }
> > }
> >
> > -static void set_filter_param(FilterParam *fp, int msize_x, int msize_y,
> float amount)
> > +static int apply_unsharp_c(AVFilterContext *ctx, AVFrame *in, AVFrame
> *out)
>
> void since the return value is ignored.
> Also some constness may help:
>
> static void apply_unsharp_c(AVFilterContext *ctx, const AVFrame *in,
> AVFrame *out)
>
> I want use the function pointer for both C code and opencl code, the
function of opencl returns int value because of there maybe some errors if
use opencl APIs.
>
> > +static void set_filter_param(UnsharpFilterParam *fp, int msize_x, int
> msize_y, float amount)
>
>
> please update docs for unsharp in doc/filters.texi
>
> Looks overall quite good. Do you have any benchmark to show?
>
Compn told me how to create the benchmark before, but sorry that I forgot
it. Can you tell me again? I will record it, thanks.
> --
> FFmpeg = Fascinating & Furious Magic Philosophical Everlasting Guru
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel at ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
More information about the ffmpeg-devel
mailing list