[FFmpeg-devel] [PATCH] libavutil/libavfilter:add opencl warpper and opencl deshake filter

Wei Gao highgod0401 at gmail.com
Sat Mar 9 13:47:29 CET 2013


Hi

Stefano, thank you for your reply, I am fixing the patch according to your
comments. I still have some questions and explanations as follows.

Thanks
Best regards

2013/3/9 Stefano Sabatini <stefasab at gmail.com>

> On date Friday 2013-03-08 10:29:49 +0800, Wei Gao encoded:
> > Hi,
> >
> > Thank you for your comments, I have some questions as follow.
> >
> > Thanks
> > Best regards
> >
> > 2013/3/8 Stefano Sabatini <stefasab at gmail.com>
> >
> > > Sorry again for the slow reply.
> > >
> > > On date Friday 2013-02-22 15:31:09 +0800, Wei Gao encoded:
> > > >
> > > > From 4824a3cb8108afe95756f46cfe5a185d3a6308dd Mon Sep 17 00:00:00
> 2001
> > > > From: highgod0401 <highgod0401 at gmail.com>
> > > > Date: Fri, 22 Feb 2013 15:25:48 +0800
> > > > Subject: [PATCH] add opencl warpper and opencl deshake filter
> > > >
> > > > ---
> > > >  configure                        |   5 +
> > > >  libavfilter/Makefile             |   3 +
> > > >  libavfilter/all_filter_kernels.c |  18 +
> > > >  libavfilter/all_filter_kernels.h |   9 +
> > > >  libavfilter/allfilters.c         |   8 +-
> > > >  libavfilter/deshake_kernel.h     | 182 +++++++++
> > > >  libavfilter/transform_opencl.c   | 157 ++++++++
> > > >  libavfilter/transform_opencl.h   |  38 ++
> > > >  libavfilter/vf_deshake.c         | 213 ++++++++++-
> > >
> > > >  libavutil/Makefile               |   4 +
> > > >  libavutil/openclwrapper.c        | 808
> > > +++++++++++++++++++++++++++++++++++++++
> > > >  libavutil/openclwrapper.h        | 198 ++++++++++
> [...]
>

In pervious mail, you comment is that openclwrapper should move to
libavfilter. but our plan is that we will submit opendecode and openencode,
these are in libavcodec, they will use openclwrapper too. And the next
OpenCL filter we plan to submit is scale, and we will add some files to
swscale folder, it will use openclwrapper too.So if we submit the patch, we
may still remove the openclwrapper to libavutil.

> > > > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
> > > > index 47158f9..e92cf8b 100644
> > > > --- a/libavfilter/allfilters.c
> > > > +++ b/libavfilter/allfilters.c
> > > > @@ -21,7 +21,9 @@
> > > >
> > > >  #include "avfilter.h"
> > > >  #include "config.h"
> > > > -
> > > > +#if CONFIG_OPENCL
> > > > +#include "all_filter_kernels.h"
> > > > +#endif
> > >
> > > avoid indirections, just insert the code here.
> > >
>
> > The function " ff_opencl_regist_kernel" is defined in openclwrapper, and
> > only can be compiled if config CONFIG_OPENCL,so should I code here and
> > still use indirections?
>
> Yes.
>
> [...]
> > > > diff --git a/libavfilter/deshake_kernel.h
> b/libavfilter/deshake_kernel.h
> > > > new file mode 100644
> [...]
> > > > +
> > > > +kernel void avfilter_transform(global  unsigned char *src,
> > > > +                               global  unsigned char *dst,
> > > > +                               global          float *matrix,
> > > > +                               global          float *matrix2,
> > >
> > > > +                                                 int interpolate,
> > > > +                                                 int fillmethod,
> > >
> > > you may use some enum here, they greatly help debugging and
> > > readability
> > >
> > Sorry, I don't get what you mean, it is a OpenCL kernel function,and each
> > parameter is set from c code using AV_OPENCLCHECK_SET_KERNEL_ARG, and
> they
> > should be set one by one.So I think it may be easy to write like this.
>
> OK if you can't import the enum definitions from transform.h.
>
> [...]
> > > > diff --git a/libavfilter/transform_opencl.c
> > > b/libavfilter/transform_opencl.c
> > > > new file mode 100644
> > > > index 0000000..4acec2b
> > > > --- /dev/null
> > > > +++ b/libavfilter/transform_opencl.c
> > > > @@ -0,0 +1,157 @@
> > > > +/*
> > > > + * Copyright (C) 2013 Wei Gao <weigao at multicorewareinc.com>
> > > > + *
> > > > + *
> > >
> > > Nit: an empty line is enough
> > >
> > > > + * 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
> > > > + */
> > > > +
> > > > +/**
> > > > + * @file
> > > > + * transform input video
> > > > + */
> > > > +
> > > > +#include "libavutil/common.h"
> > > > +#include "libavutil/avassert.h"
> > > > +#include "libavutil/avstring.h"
> > > > +#include "libavutil/openclwrapper.h"
> > > > +#include "transform.h"
> > > > +#include "transform_opencl.h"
> > > > +
> > > > +
> > > > +
> > > > +static int ff_filter_transform_func(void **userdata, AV_KernelEnv
> *kenv)
> > > > +{
> > > > +    cl_mem src = (cl_mem)userdata[0];
> > > > +    cl_mem dst = (cl_mem)userdata[1];
> > > > +    int src_stride_lu = (int)userdata[2];
> > > > +    int dst_stride_lu = (int)userdata[3];
> > > > +    int src_stride_ch = (int)userdata[4];
> > > > +    int dst_stride_ch = (int)userdata[5];
> > > > +    int width      = (int)userdata[6];
> > > > +    int height     = (int)userdata[7];
> > > > +    int cw          = (int)userdata[8];
> > > > +    int ch     = (int)userdata[9];
> > > > +    float *matrix  = (float *)userdata[10];
> > > > +    float *matrix2  = (float *)userdata[11];
> > > > +    int interpolate = (int)userdata[12];
> > > > +    int fillmethod  = (int)userdata[13];
> > > > +    cl_mem matrix_buf = (cl_mem)userdata[14];
> > > > +    cl_mem matrix_buf2  = (cl_mem)userdata[15];
> > > > +    AV_KernelEnv *env  = (AV_KernelEnv *)userdata[16];
> > > > +    cl_uint status;
> > > > +    void *mapped;
> > > > +    const size_t global_work_size = width * height + 2 * ch * cw;
> > > > +    int m_size = 6;
> > > > +    cl_kernel kernel;
> > > > +    int arg_no;
> > > > +
> > > > +
> > > > +    mapped = clEnqueueMapBuffer(kenv->command_queue, matrix_buf,
> > > CL_TRUE, CL_MAP_WRITE, 0, m_size*sizeof(cl_float),0,NULL, NULL, NULL);
> > > > +    memcpy(mapped,matrix,m_size*sizeof(cl_float));
> > > > +    clEnqueueUnmapMemObject(kenv->command_queue, matrix_buf,
> mapped, 0,
> > > NULL, NULL);
> > > > +
> > > > +    mapped = clEnqueueMapBuffer(kenv->command_queue, matrix_buf2,
> > > CL_TRUE, CL_MAP_WRITE, 0, m_size*sizeof(cl_float),0,NULL, NULL, NULL);
> > > > +    memcpy(mapped,matrix2,m_size*sizeof(cl_float));
> > > > +    clEnqueueUnmapMemObject(kenv->command_queue, matrix_buf2,
> mapped,
> > > 0, NULL, NULL);
> > > > +
> > > > +    if (!env->kernel) {
> > > > +        status =  av_opencl_create_kernel("avfilter_transform",
> kenv);
> > > > +        if (status) {
> > > > +            av_log(NULL,AV_LOG_ERROR,"clCreateKernel Error
> > > %s\n","avfilter_transform");
> > > > +            return 0;
> > > > +        }
> > > > +        env->command_queue = kenv->command_queue;
> > > > +        env->context = kenv->context;
> > > > +        env->kernel = kenv->kernel;
> > > > +        av_strlcpy(env->kernel_name,kenv->kernel_name,150);
> > > > +        env->program = kenv->program;
> > > > +    }
> > > > +    kernel = env->kernel;
> > > > +    arg_no = 0;
> > >
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&src);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&dst);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&matrix_buf);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_mem),&matrix_buf2);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&interpolate);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&fillmethod);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&src_stride_lu);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&dst_stride_lu);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&src_stride_ch);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&dst_stride_ch);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&height);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&width);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&ch);
> > > > +    AV_OPENCLCHECK_SET_KERNEL_ARG(sizeof(cl_int),&cw);
> > > > +
> > > > +    AV_OPENCLCHECK( clEnqueueNDRangeKernel, env->command_queue,
> > > env->kernel, 1, NULL,
> > > > +              &global_work_size, NULL, 0, NULL, NULL);
> > >
> > > since this is the only user, move the macro definition locally to this
> > > file (rule: avoid indirections unless strictly necessary).
> > >
> >
>
> >  This is because we only submit deshake filter recently,other OpenCL
> > filters we will submit later will use the macro
> > AV_OPENCLCHECK_SET_KERNEL_ARG and AV_OPENCLCHECK to easy the code
> write.So
> > we add it in opencl wraper.If not, every OpenCL filter will add the two
> > macro.
>
> OK, but please use names:
>
> AV_OPENCL_SET_KERNEL_ARG
> AV_OPENCL_CHECK
>
> also you could probably move sizeof() within the macro.
>
> [...]
> > > > diff --git a/libavfilter/transform_opencl.h
> > > b/libavfilter/transform_opencl.h
> > > > new file mode 100644
> > > > index 0000000..aa1b4e0
> > > > --- /dev/null
> > > > +++ b/libavfilter/transform_opencl.h
> > > > @@ -0,0 +1,38 @@
> > > > +/*
> > > > + * 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_TRANSFORM_OPENCL_H
> > > > +#define AVFILTER_TRANSFORM_OPENCL_H
> > > > +
> > > > +#include <stdint.h>
> > > > +
> > > > +void avfilter_transform_cl( void *src,  void *dst,
> > > > +                        int src_stride_lu, int dst_stride_lu,
> > > > +                        int src_stride_ch, int dst_stride_ch,
> > > > +                        int width, int height, int cw, int ch,
> > > > +                        const float *matrix, const float *matrix2,
> > > > +                        const void *matrix_cl, const void
> *matrix2_cl,
> > >
> > > > +                        enum InterpolateMethod interpolate,
> > > > +                        enum FillMethod fill, AV_KernelEnv *env);
> > >
> > > where are these enums defined?
> > >
> > They are defined in "transform.h", I move the header file here.
> >
> > >
> > > > +int ff_init_transform(void);
> > >
> > > Is this file supposed to be public? If not you should avoid to use
> > > avfilter_ prefix (reserved for public functions).
> > >
> > > I suggest something along the lines of:
> > >
> > > ff_opencl_transform()
> > > ff_opencl_transform_init()
> > >
>
> > This file,I referenced "transform.h",this file will be used in
> > deshake_opencl filter.
>
> Thus it shouldn't need to be made public, and the ff_prefix should be
> used.
>
> [...]
> > > > +AVFilter avfilter_vf_deshake_opencl = {
> > > > +    .name          = "deshake_opencl",
> > > > +    .description   = NULL_IF_CONFIG_SMALL("Stabilize shaky video
> using
> > > OpenCL."),
> > > > +    .priv_size     = sizeof(DeshakeContext),
> > > > +    .init          = init_opencl,
> > > > +    .uninit        = uninit_opencl,
> > > > +    .query_formats = query_formats,
> > > > +    .inputs        = deshake_opencl_inputs,
> > > > +    .outputs       = deshake_opencl_outputs,
> > > > +};
> > > > +#endif
> > >
> > > General question: is this filter supposed to be binary compatible with
> > > deshake? Or it just shares the same syntax?
> > >
> >
> > Sorry, I don't get what your mean.Our idea is to make the deshake_opencl
> > filter can be used both in ffmpeg.exe and a filter in libavfilter.lib
> witch
> > can be used in other application program.
>
> Is the output created by the opencl deshake filter the same as the one
> produced by deshake?
>
> You can test using:
> ffmpeg -f lavfi -i testsrc=d=10,deshake -f md5 - -nostats
> ffmpeg -f lavfi -i testsrc=d=10,deshake_opencl -f md5 - -nostats
>
> and comparing the output, or using the showinfo filter.
>
> [...]
> > > > +int av_opencl_create_kernel(const char * kernelname, AV_KernelEnv *
> env)
> > > > +{
> > > > +    int status;
> > > > +    env->kernel = clCreateKernel(gpu_env.programs[0], kernelname,
> > > &status);
> > > > +    env->context = gpu_env.context;
> > > > +    env->command_queue = gpu_env.command_queue;
> > >
> > > > +    if (status != CL_SUCCESS) {
> > > > +        av_log(&openclwrapperutils,AV_LOG_ERROR,"status =
> %d\n",status);
> > >
> > > Error messages should be intelligible, reporting just a status is not
> > > enough.
> > >
> >
>
> > There are a lot of statuses,and you can find the status in CL.h(this file
> > is the OpenCL SDK),so I think just return the status then you can find
> the
> > error.
>
> > Also if you think it should return the error exactly as string, I
> > will write a function to get the error string.
>
> Yes. The point is that software should do its best to explain what's
> going wrong. Returning a number is not particularly useful to the
> user, a string containing an error description should be already better:
>
> Something like this should be good:
> if (status != CL_SUCCESS) {
>     av_log(&openclwrapperutils, AV_LOG_ERROR, "Error creating OpenCL
> kernel: %s", opencl_errstr(status));
>
> [...]
> > >
> (av_opencl_read_cl_buffer(cl_inBuf,gpu_env.temp_buffer,sizeof(uint8_t)*(linesize0
> > > + linesize1)*height)) {
> > > > +        memcpy(Ybuf,gpu_env.temp_buffer,linesize0 * height);
> > > > +        memcpy(Ubuf,gpu_env.temp_buffer + linesize0 *
> height,linesize1
> > > *chrH);
> > > > +        memcpy(Vbuf,gpu_env.temp_buffer + linesize0 * height +
> > > linesize1 * chrH,linesize2 * chrH);
> > > > +    }
> > > > +    return 1;
> > > > +}
> > > > +void av_opencl_save_buffer(const char *filtername,void *cl_inbuf)
> > > > +{
> > > > +    int i = 0;
> > > > +    while (strlen(filter_buffer[i].filter_name)) {
> > > > +        i++;
> > > > +    }
> > > > +    if (i > (MAX_FILTER_NUM - 1)) {
> > > > +        av_log(&openclwrapperutils,AV_LOG_ERROR,"filter num is too
> > > large\n");
> > > > +        return;
> > > > +    }
> > > > +    if(strlen(filtername) > MAX_FILTER_NAME_LEN) {
> > > > +        av_log(&openclwrapperutils,AV_LOG_ERROR,"filter name is too
> > > long\n");
> > > > +        return;
> > > > +    }
> > > > +
> > >
>  av_strlcpy(filter_buffer[i].filter_name,filtername,MAX_FILTER_NAME_LEN+1);
> > > > +    filter_buffer[i].cl_inbuf = cl_inbuf;
> > > > +    return;
> > >
> > > aren't you supposed to return an error in case of too long buffers?
> > >
> >
>
> > This functions is used to save the OpenCL memory pointer and if the
> follow
> > filter is OpenCL filter, it can get the OpenCL buffer as input, it can
> > avoid copy data between GPU and CPU.So the "cl_inbuf" is the buffer
> > per-OpenCL filter has created,this function just save the pointer.
>
> Yes but what happens if the buffer are too short to contain the copied
> data?
>
Our design is that the buffer will be used as the input parameter to
filter.and it is read only,so the filter will not copy the data to buffer.

>
> [...]
> > > > diff --git a/libavutil/openclwrapper.h b/libavutil/openclwrapper.h
>
> Keep the name short, opencl.h should be enough (assuming we are going
> to keep this file in libavutil).
>
> > > > new file mode 100644
> > > > index 0000000..a3ac8c2
> > > > --- /dev/null
> > > > +++ b/libavutil/openclwrapper.h
> > > > @@ -0,0 +1,198 @@
> > > > +/*
> > > > + * 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 "config.h"
> > > > +
> > > > +#ifndef LIBAVUTIL_OPENCLWRAPPER_H
> > > > +#define LIBAVUTIL_OPENCLWRAPPER_H
> > > > +
> > > > +#include <CL/cl.h>
> > > > +
> > >
> > > > +#define AV_OPENCLCHECK(method, ...)\
> > > > +    status = method( __VA_ARGS__ ); if( status != CL_SUCCESS ) {\
> > > > +        av_log(NULL,AV_LOG_ERROR, " error %s %d\n", # method,
> status );
> > >  return status; }
> > > > +
> > > > +#define AV_OPENCLCHECK_SET_KERNEL_ARG(arg_size,arg_ptr)
> AV_OPENCLCHECK(
> > > clSetKernelArg,kernel,(arg_no++),(arg_size),(void*)(arg_ptr))
> > >
> > > avoid nested macros
> > >
> > > > +
> > > > +#define FF_OPENCL_KERNEL( ... )# __VA_ARGS__
> > > > +
> > > > +typedef struct AV_KernelEnv {
> > > > +    cl_context context;
> > > > +    cl_command_queue command_queue;
> > > > +    cl_program program;
> > > > +    cl_kernel kernel;
> > > > +    char kernel_name[150];
> > > > +}AV_KernelEnv;
> > >
> > > AVKernelEnv?
> > >
> > > > +
> > > > +typedef struct AV_ExtOpenCLInfo {
> > > > +    cl_platform_id platform;
> > >
> > > > +    cl_device_type devide_type;
> > >
> > > typo?
> > >
> > > > +    cl_context context;
> > > > +    cl_device_id *devices;
> > >
> > > devices_id?
> > >
> > > > +    cl_device_id  dev;
> > > > +    cl_command_queue command_queue;
> > > > +    char *platformName;
> > >
> > > camelCase
> > >
> > > > +}AV_ExtOpenCLInfo;
> > >
> > > AVExtOpenCLInfo
> > >
> > > also what's "Ext" for?
> > >
> >
>
> > Let me explain it. There are two cases to use OpenCL filters, the first
> one
> > is just as ffmpeg.exe,the application program doesn't create OpenCL
> >  environment but it want to use OpenCL filter in ffmpeg or
> > libavfilter.lib,under this condition,libavfilter(ffmpeg) should create
> the
> > OpenCL filter it self,so it doesn't use AVExtOpenCLInfo.The other one is
> > that the application program has created the OpenCL environment, under
> this
> > condition, libavfilter(ffmpeg) may use the OpenCL environment wtich
> > application program created,so the application program just set the
> > AVExtOpenCLInfo and pass it two the API "av_opencl_init_run_env",
> > libavfilter will use it to init opencl filter.
>
> Why two distinct interfaces if just one will be used?
>
> Also my comment was more about what "Ext" stands for, I suppose it
> stands for "External". In each case the name should be something like
> AVOpenCLExtInfo. The usual convention here is library prefix ("AV"),
> followed by module name ("OpenCL"), followed by a possibly descriptive
> name.
>
This will be as an input of  "av_opencl_init_run_env" if the application
has init the opencl. And if the app use the filter without init OpenCL, it
will instead by NULL. Sorry that I spell wrong  "it two the API", it should
be " it to the API".

>
> [...]
> --
> FFmpeg = Fierce & Faithless Mean Power Enlightening 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